Merge Mesa 21.3.7

10 files changed, 237 insertions, 389 deletions

diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_build.c b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_build.c
index 9e1a55342..73d1eb26a 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_build.c
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_build.c
@@ -218,7 +218,6 @@ tgsi_default_declaration_interp( void )
 
    di.Interpolate = TGSI_INTERPOLATE_CONSTANT;
    di.Location = TGSI_INTERPOLATE_LOC_CENTER;
-   di.CylindricalWrap = 0;
    di.Padding = 0;
 
    return di;
@@ -227,7 +226,6 @@ tgsi_default_declaration_interp( void )
 static struct tgsi_declaration_interp
 tgsi_build_declaration_interp(unsigned interpolate,
                               unsigned interpolate_location,
-                              unsigned cylindrical_wrap,
                               struct tgsi_declaration *declaration,
                               struct tgsi_header *header)
 {
@@ -235,7 +233,6 @@ tgsi_build_declaration_interp(unsigned interpolate,
 
    di.Interpolate = interpolate;
    di.Location = interpolate_location;
-   di.CylindricalWrap = cylindrical_wrap;
    di.Padding = 0;
 
    declaration_grow(declaration, header);
@@ -467,7 +464,6 @@ tgsi_build_full_declaration(
 
       *di = tgsi_build_declaration_interp(full_decl->Interp.Interpolate,
                                           full_decl->Interp.Location,
-                                          full_decl->Interp.CylindricalWrap,
                                           declaration,
                                           header);
    }
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_dump.c b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_dump.c
index 110aeb608..68e30b6b8 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_dump.c
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_dump.c
@@ -436,22 +436,6 @@ iter_declaration(
          TXT( ", " );
          ENM( decl->Interp.Location, tgsi_interpolate_locations );
       }
-
-      if (decl->Interp.CylindricalWrap) {
-         TXT(", CYLWRAP_");
-         if (decl->Interp.CylindricalWrap & TGSI_CYLINDRICAL_WRAP_X) {
-            CHR('X');
-         }
-         if (decl->Interp.CylindricalWrap & TGSI_CYLINDRICAL_WRAP_Y) {
-            CHR('Y');
-         }
-         if (decl->Interp.CylindricalWrap & TGSI_CYLINDRICAL_WRAP_Z) {
-            CHR('Z');
-         }
-         if (decl->Interp.CylindricalWrap & TGSI_CYLINDRICAL_WRAP_W) {
-            CHR('W');
-         }
-      }
    }
 
    if (decl->Declaration.Invariant) {
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_exec.c b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_exec.c
index fbfb04ef3..a8446ff27 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_exec.c
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_exec.c
@@ -68,8 +68,6 @@
 #define DEBUG_EXECUTION 0
 
 
-#define FAST_MATH 0
-
 #define TILE_TOP_LEFT     0
 #define TILE_TOP_RIGHT    1
 #define TILE_BOTTOM_LEFT  2
@@ -80,9 +78,9 @@ union tgsi_double_channel {
    unsigned u[TGSI_QUAD_SIZE][2];
    uint64_t u64[TGSI_QUAD_SIZE];
    int64_t i64[TGSI_QUAD_SIZE];
-};
+} ALIGN16;
 
-struct tgsi_double_vector {
+struct ALIGN16 tgsi_double_vector {
    union tgsi_double_channel xy;
    union tgsi_double_channel zw;
 };
@@ -263,20 +261,20 @@ static void
 micro_dmax(union tgsi_double_channel *dst,
            const union tgsi_double_channel *src)
 {
-   dst->d[0] = src[0].d[0] > src[1].d[0] || isnan(src[1].d[0]) ? src[0].d[0] : src[1].d[0];
-   dst->d[1] = src[0].d[1] > src[1].d[1] || isnan(src[1].d[1]) ? src[0].d[1] : src[1].d[1];
-   dst->d[2] = src[0].d[2] > src[1].d[2] || isnan(src[1].d[2]) ? src[0].d[2] : src[1].d[2];
-   dst->d[3] = src[0].d[3] > src[1].d[3] || isnan(src[1].d[3]) ? src[0].d[3] : src[1].d[3];
+   dst->d[0] = fmax(src[0].d[0], src[1].d[0]);
+   dst->d[1] = fmax(src[0].d[1], src[1].d[1]);
+   dst->d[2] = fmax(src[0].d[2], src[1].d[2]);
+   dst->d[3] = fmax(src[0].d[3], src[1].d[3]);
 }
 
 static void
 micro_dmin(union tgsi_double_channel *dst,
            const union tgsi_double_channel *src)
 {
-   dst->d[0] = src[0].d[0] < src[1].d[0] || isnan(src[1].d[0]) ? src[0].d[0] : src[1].d[0];
-   dst->d[1] = src[0].d[1] < src[1].d[1] || isnan(src[1].d[1]) ? src[0].d[1] : src[1].d[1];
-   dst->d[2] = src[0].d[2] < src[1].d[2] || isnan(src[1].d[2]) ? src[0].d[2] : src[1].d[2];
-   dst->d[3] = src[0].d[3] < src[1].d[3] || isnan(src[1].d[3]) ? src[0].d[3] : src[1].d[3];
+   dst->d[0] = fmin(src[0].d[0], src[1].d[0]);
+   dst->d[1] = fmin(src[0].d[1], src[1].d[1]);
+   dst->d[2] = fmin(src[0].d[2], src[1].d[2]);
+   dst->d[3] = fmin(src[0].d[3], src[1].d[3]);
 }
 
 static void
@@ -415,12 +413,6 @@ static void
 micro_exp2(union tgsi_exec_channel *dst,
            const union tgsi_exec_channel *src)
 {
-#if FAST_MATH
-   dst->f[0] = util_fast_exp2(src->f[0]);
-   dst->f[1] = util_fast_exp2(src->f[1]);
-   dst->f[2] = util_fast_exp2(src->f[2]);
-   dst->f[3] = util_fast_exp2(src->f[3]);
-#else
 #if DEBUG
    /* Inf is okay for this instruction, so clamp it to silence assertions. */
    uint i;
@@ -442,7 +434,6 @@ micro_exp2(union tgsi_exec_channel *dst,
    dst->f[1] = powf(2.0f, src->f[1]);
    dst->f[2] = powf(2.0f, src->f[2]);
    dst->f[3] = powf(2.0f, src->f[3]);
-#endif /* FAST_MATH */
 }
 
 static void
@@ -509,17 +500,10 @@ static void
 micro_lg2(union tgsi_exec_channel *dst,
           const union tgsi_exec_channel *src)
 {
-#if FAST_MATH
-   dst->f[0] = util_fast_log2(src->f[0]);
-   dst->f[1] = util_fast_log2(src->f[1]);
-   dst->f[2] = util_fast_log2(src->f[2]);
-   dst->f[3] = util_fast_log2(src->f[3]);
-#else
    dst->f[0] = logf(src->f[0]) * 1.442695f;
    dst->f[1] = logf(src->f[1]) * 1.442695f;
    dst->f[2] = logf(src->f[2]) * 1.442695f;
    dst->f[3] = logf(src->f[3]) * 1.442695f;
-#endif
 }
 
 static void
@@ -1075,9 +1059,6 @@ tgsi_exec_machine_bind_shader(
    tgsi_dump(tokens, 0);
 #endif
 
-   util_init_math();
-
-
    mach->Tokens = tokens;
    mach->Sampler = sampler;
    mach->Image = image;
@@ -1357,10 +1338,10 @@ micro_max(union tgsi_exec_channel *dst,
           const union tgsi_exec_channel *src0,
           const union tgsi_exec_channel *src1)
 {
-   dst->f[0] = src0->f[0] > src1->f[0] || isnan(src1->f[0]) ? src0->f[0] : src1->f[0];
-   dst->f[1] = src0->f[1] > src1->f[1] || isnan(src1->f[1]) ? src0->f[1] : src1->f[1];
-   dst->f[2] = src0->f[2] > src1->f[2] || isnan(src1->f[2]) ? src0->f[2] : src1->f[2];
-   dst->f[3] = src0->f[3] > src1->f[3] || isnan(src1->f[3]) ? src0->f[3] : src1->f[3];
+   dst->f[0] = fmaxf(src0->f[0], src1->f[0]);
+   dst->f[1] = fmaxf(src0->f[1], src1->f[1]);
+   dst->f[2] = fmaxf(src0->f[2], src1->f[2]);
+   dst->f[3] = fmaxf(src0->f[3], src1->f[3]);
 }
 
 static void
@@ -1368,10 +1349,10 @@ micro_min(union tgsi_exec_channel *dst,
           const union tgsi_exec_channel *src0,
           const union tgsi_exec_channel *src1)
 {
-   dst->f[0] = src0->f[0] < src1->f[0] || isnan(src1->f[0]) ? src0->f[0] : src1->f[0];
-   dst->f[1] = src0->f[1] < src1->f[1] || isnan(src1->f[1]) ? src0->f[1] : src1->f[1];
-   dst->f[2] = src0->f[2] < src1->f[2] || isnan(src1->f[2]) ? src0->f[2] : src1->f[2];
-   dst->f[3] = src0->f[3] < src1->f[3] || isnan(src1->f[3]) ? src0->f[3] : src1->f[3];
+   dst->f[0] = fminf(src0->f[0], src1->f[0]);
+   dst->f[1] = fminf(src0->f[1], src1->f[1]);
+   dst->f[2] = fminf(src0->f[2], src1->f[2]);
+   dst->f[3] = fminf(src0->f[3], src1->f[3]);
 }
 
 static void
@@ -1402,17 +1383,10 @@ micro_pow(
    const union tgsi_exec_channel *src0,
    const union tgsi_exec_channel *src1 )
 {
-#if FAST_MATH
-   dst->f[0] = util_fast_pow( src0->f[0], src1->f[0] );
-   dst->f[1] = util_fast_pow( src0->f[1], src1->f[1] );
-   dst->f[2] = util_fast_pow( src0->f[2], src1->f[2] );
-   dst->f[3] = util_fast_pow( src0->f[3], src1->f[3] );
-#else
    dst->f[0] = powf( src0->f[0], src1->f[0] );
    dst->f[1] = powf( src0->f[1], src1->f[1] );
    dst->f[2] = powf( src0->f[2], src1->f[2] );
    dst->f[3] = powf( src0->f[3], src1->f[3] );
-#endif
 }
 
 static void
@@ -1452,28 +1426,22 @@ fetch_src_file_channel(const struct tgsi_exec_machine *mach,
    switch (file) {
    case TGSI_FILE_CONSTANT:
       for (i = 0; i < TGSI_QUAD_SIZE; i++) {
-         assert(index2D->i[i] >= 0 && index2D->i[i] < PIPE_MAX_CONSTANT_BUFFERS);
-
-         if (index->i[i] < 0) {
+         /* NOTE: copying the const value as a uint instead of float */
+         const uint constbuf = index2D->i[i];
+         const unsigned pos = index->i[i] * 4 + swizzle;
+         /* const buffer bounds check */
+         if (pos >= mach->ConstsSize[constbuf] / 4) {
+            if (0) {
+               /* Debug: print warning */
+               static int count = 0;
+               if (count++ < 100)
+                  debug_printf("TGSI Exec: const buffer index %d"
+                                 " out of bounds\n", pos);
+            }
             chan->u[i] = 0;
          } else {
-            /* NOTE: copying the const value as a uint instead of float */
-            const uint constbuf = index2D->i[i];
-            const int pos = index->i[i] * 4 + swizzle;
-            /* const buffer bounds check */
-            if (pos < 0 || pos >= (int) mach->ConstsSize[constbuf] / 4) {
-               if (0) {
-                  /* Debug: print warning */
-                  static int count = 0;
-                  if (count++ < 100)
-                     debug_printf("TGSI Exec: const buffer index %d"
-                                  " out of bounds\n", pos);
-               }
-               chan->u[i] = 0;
-            } else {
-               const uint *buf = (const uint *)mach->Consts[constbuf];
-               chan->u[i] = buf[pos];
-            }
+            const uint *buf = (const uint *)mach->Consts[constbuf];
+            chan->u[i] = buf[pos];
          }
       }
       break;
@@ -1730,19 +1698,13 @@ static union tgsi_exec_channel *
 store_dest_dstret(struct tgsi_exec_machine *mach,
                  const union tgsi_exec_channel *chan,
                  const struct tgsi_full_dst_register *reg,
-                 uint chan_index,
-                 enum tgsi_exec_datatype dst_datatype)
+                 uint chan_index)
 {
    static union tgsi_exec_channel null;
    union tgsi_exec_channel *dst;
-   union tgsi_exec_channel index2D;
    int offset = 0;  /* indirection offset */
    int index;
 
-   /* for debugging */
-   if (0 && dst_datatype == TGSI_EXEC_DATA_FLOAT) {
-      check_inf_or_nan(chan);
-   }
 
    /* There is an extra source register that indirectly subscripts
     * a register file. The direct index now becomes an offset
@@ -1780,77 +1742,6 @@ store_dest_dstret(struct tgsi_exec_machine *mach,
       offset = indir_index.i[0];
    }
 
-   /* There is an extra source register that is a second
-    * subscript to a register file. Effectively it means that
-    * the register file is actually a 2D array of registers.
-    *
-    *    file[3][1],
-    *    where:
-    *       [3] = Dimension.Index
-    */
-   if (reg->Register.Dimension) {
-      index2D.i[0] =
-      index2D.i[1] =
-      index2D.i[2] =
-      index2D.i[3] = reg->Dimension.Index;
-
-      /* Again, the second subscript index can be addressed indirectly
-       * identically to the first one.
-       * Nothing stops us from indirectly addressing the indirect register,
-       * but there is no need for that, so we won't exercise it.
-       *
-       *    file[ind[4].y+3][1],
-       *    where:
-       *       ind = DimIndirect.File
-       *       [4] = DimIndirect.Index
-       *       .y = DimIndirect.SwizzleX
-       */
-      if (reg->Dimension.Indirect) {
-         union tgsi_exec_channel index2;
-         union tgsi_exec_channel indir_index;
-         const uint execmask = mach->ExecMask;
-         unsigned swizzle;
-         uint i;
-
-         index2.i[0] =
-         index2.i[1] =
-         index2.i[2] =
-         index2.i[3] = reg->DimIndirect.Index;
-
-         swizzle = reg->DimIndirect.Swizzle;
-         fetch_src_file_channel(mach,
-                                reg->DimIndirect.File,
-                                swizzle,
-                                &index2,
-                                &ZeroVec,
-                                &indir_index);
-
-         index2D.i[0] += indir_index.i[0];
-         index2D.i[1] += indir_index.i[1];
-         index2D.i[2] += indir_index.i[2];
-         index2D.i[3] += indir_index.i[3];
-
-         /* for disabled execution channels, zero-out the index to
-          * avoid using a potential garbage value.
-          */
-         for (i = 0; i < TGSI_QUAD_SIZE; i++) {
-            if ((execmask & (1 << i)) == 0) {
-               index2D.i[i] = 0;
-            }
-         }
-      }
-
-      /* If by any chance there was a need for a 3D array of register
-       * files, we would have to check whether Dimension is followed
-       * by a dimension register and continue the saga.
-       */
-   } else {
-      index2D.i[0] =
-      index2D.i[1] =
-      index2D.i[2] =
-      index2D.i[3] = 0;
-   }
-
    switch (reg->Register.File) {
    case TGSI_FILE_NULL:
       dst = &null;
@@ -1886,8 +1777,7 @@ store_dest_dstret(struct tgsi_exec_machine *mach,
       break;
 
    default:
-      assert( 0 );
-      return NULL;
+      unreachable("Bad destination file");
    }
 
    return dst;
@@ -1897,14 +1787,13 @@ static void
 store_dest_double(struct tgsi_exec_machine *mach,
                  const union tgsi_exec_channel *chan,
                  const struct tgsi_full_dst_register *reg,
-                 uint chan_index,
-                 enum tgsi_exec_datatype dst_datatype)
+                 uint chan_index)
 {
    union tgsi_exec_channel *dst;
    const uint execmask = mach->ExecMask;
    int i;
 
-   dst = store_dest_dstret(mach, chan, reg, chan_index, dst_datatype);
+   dst = store_dest_dstret(mach, chan, reg, chan_index);
    if (!dst)
       return;
 
@@ -1919,14 +1808,13 @@ store_dest(struct tgsi_exec_machine *mach,
            const union tgsi_exec_channel *chan,
            const struct tgsi_full_dst_register *reg,
            const struct tgsi_full_instruction *inst,
-           uint chan_index,
-           enum tgsi_exec_datatype dst_datatype)
+           uint chan_index)
 {
    union tgsi_exec_channel *dst;
    const uint execmask = mach->ExecMask;
    int i;
 
-   dst = store_dest_dstret(mach, chan, reg, chan_index, dst_datatype);
+   dst = store_dest_dstret(mach, chan, reg, chan_index);
    if (!dst)
       return;
 
@@ -1937,14 +1825,8 @@ store_dest(struct tgsi_exec_machine *mach,
    }
    else {
       for (i = 0; i < TGSI_QUAD_SIZE; i++)
-         if (execmask & (1 << i)) {
-            if (chan->f[i] < 0.0f || isnan(chan->f[i]))
-               dst->f[i] = 0.0f;
-            else if (chan->f[i] > 1.0f)
-               dst->f[i] = 1.0f;
-            else
-               dst->i[i] = chan->i[i];
-         }
+         if (execmask & (1 << i))
+            dst->f[i] = fminf(fmaxf(chan->f[i], 0.0f), 1.0f);
    }
 }
 
@@ -2311,7 +2193,7 @@ exec_tex(struct tgsi_exec_machine *mach,
 
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -2355,12 +2237,10 @@ exec_lodq(struct tgsi_exec_machine *mach,
                             r[1].f);
 
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_X) {
-      store_dest(mach, &r[0], &inst->Dst[0], inst, TGSI_CHAN_X,
-                 TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r[0], &inst->Dst[0], inst, TGSI_CHAN_X);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Y) {
-      store_dest(mach, &r[1], &inst->Dst[0], inst, TGSI_CHAN_Y,
-                 TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r[1], &inst->Dst[0], inst, TGSI_CHAN_Y);
    }
    if (inst->Instruction.Opcode == TGSI_OPCODE_LOD) {
       unsigned char swizzles[4];
@@ -2374,21 +2254,19 @@ exec_lodq(struct tgsi_exec_machine *mach,
          if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
             if (swizzles[chan] >= 2) {
                store_dest(mach, &ZeroVec,
-                          &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+                          &inst->Dst[0], inst, chan);
             } else {
                store_dest(mach, &r[swizzles[chan]],
-                          &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+                          &inst->Dst[0], inst, chan);
             }
          }
       }
    } else {
       if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_X) {
-         store_dest(mach, &r[0], &inst->Dst[0], inst, TGSI_CHAN_X,
-                    TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &r[0], &inst->Dst[0], inst, TGSI_CHAN_X);
       }
       if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Y) {
-         store_dest(mach, &r[1], &inst->Dst[0], inst, TGSI_CHAN_Y,
-                    TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &r[1], &inst->Dst[0], inst, TGSI_CHAN_Y);
       }
    }
 }
@@ -2495,7 +2373,7 @@ exec_txd(struct tgsi_exec_machine *mach,
 
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -2573,14 +2451,14 @@ exec_txf(struct tgsi_exec_machine *mach,
       for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
          if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
             store_dest(mach, &r[swizzles[chan]],
-                       &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+                       &inst->Dst[0], inst, chan);
          }
       }
    }
    else {
       for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
          if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-            store_dest(mach, &r[chan], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+            store_dest(mach, &r[chan], &inst->Dst[0], inst, chan);
          }
       }
    }
@@ -2611,8 +2489,7 @@ exec_txq(struct tgsi_exec_machine *mach,
 
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan,
-                    TGSI_EXEC_DATA_INT);
+         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -2746,7 +2623,7 @@ exec_sample(struct tgsi_exec_machine *mach,
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
          store_dest(mach, &r[swizzles[chan]],
-                    &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+                    &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -2828,7 +2705,7 @@ exec_sample_d(struct tgsi_exec_machine *mach,
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
          store_dest(mach, &r[swizzles[chan]],
-                    &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+                    &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3054,7 +2931,6 @@ static void
 exec_scalar_unary(struct tgsi_exec_machine *mach,
                   const struct tgsi_full_instruction *inst,
                   micro_unary_op op,
-                  enum tgsi_exec_datatype dst_datatype,
                   enum tgsi_exec_datatype src_datatype)
 {
    unsigned int chan;
@@ -3065,7 +2941,7 @@ exec_scalar_unary(struct tgsi_exec_machine *mach,
    op(&dst, &src);
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst, &inst->Dst[0], inst, chan, dst_datatype);
+         store_dest(mach, &dst, &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3074,7 +2950,6 @@ static void
 exec_vector_unary(struct tgsi_exec_machine *mach,
                   const struct tgsi_full_instruction *inst,
                   micro_unary_op op,
-                  enum tgsi_exec_datatype dst_datatype,
                   enum tgsi_exec_datatype src_datatype)
 {
    unsigned int chan;
@@ -3090,7 +2965,7 @@ exec_vector_unary(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan, dst_datatype);
+         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3103,7 +2978,6 @@ static void
 exec_scalar_binary(struct tgsi_exec_machine *mach,
                    const struct tgsi_full_instruction *inst,
                    micro_binary_op op,
-                   enum tgsi_exec_datatype dst_datatype,
                    enum tgsi_exec_datatype src_datatype)
 {
    unsigned int chan;
@@ -3115,7 +2989,7 @@ exec_scalar_binary(struct tgsi_exec_machine *mach,
    op(&dst, &src[0], &src[1]);
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst, &inst->Dst[0], inst, chan, dst_datatype);
+         store_dest(mach, &dst, &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3124,7 +2998,6 @@ static void
 exec_vector_binary(struct tgsi_exec_machine *mach,
                    const struct tgsi_full_instruction *inst,
                    micro_binary_op op,
-                   enum tgsi_exec_datatype dst_datatype,
                    enum tgsi_exec_datatype src_datatype)
 {
    unsigned int chan;
@@ -3141,7 +3014,7 @@ exec_vector_binary(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan, dst_datatype);
+         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3155,7 +3028,6 @@ static void
 exec_vector_trinary(struct tgsi_exec_machine *mach,
                     const struct tgsi_full_instruction *inst,
                     micro_trinary_op op,
-                    enum tgsi_exec_datatype dst_datatype,
                     enum tgsi_exec_datatype src_datatype)
 {
    unsigned int chan;
@@ -3173,7 +3045,7 @@ exec_vector_trinary(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan, dst_datatype);
+         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3188,7 +3060,6 @@ static void
 exec_vector_quaternary(struct tgsi_exec_machine *mach,
                        const struct tgsi_full_instruction *inst,
                        micro_quaternary_op op,
-                       enum tgsi_exec_datatype dst_datatype,
                        enum tgsi_exec_datatype src_datatype)
 {
    unsigned int chan;
@@ -3207,7 +3078,7 @@ exec_vector_quaternary(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan, dst_datatype);
+         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3231,7 +3102,7 @@ exec_dp3(struct tgsi_exec_machine *mach,
 
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &arg[2], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &arg[2], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3255,7 +3126,7 @@ exec_dp4(struct tgsi_exec_machine *mach,
 
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &arg[2], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &arg[2], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3277,7 +3148,7 @@ exec_dp2(struct tgsi_exec_machine *mach,
 
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &arg[2], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &arg[2], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3297,7 +3168,7 @@ exec_pk2h(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst, &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_UINT);
+         store_dest(mach, &dst, &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3316,7 +3187,7 @@ exec_up2h(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst[chan & 1], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &dst[chan & 1], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3355,8 +3226,7 @@ exec_ucmp(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan,
-                    TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &dst.xyzw[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3381,16 +3251,16 @@ exec_dst(struct tgsi_exec_machine *mach,
    }
 
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_X) {
-      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_X, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_X);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Y) {
-      store_dest(mach, &d[TGSI_CHAN_Y], &inst->Dst[0], inst, TGSI_CHAN_Y, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &d[TGSI_CHAN_Y], &inst->Dst[0], inst, TGSI_CHAN_Y);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Z) {
-      store_dest(mach, &d[TGSI_CHAN_Z], &inst->Dst[0], inst, TGSI_CHAN_Z, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &d[TGSI_CHAN_Z], &inst->Dst[0], inst, TGSI_CHAN_Z);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_W) {
-      store_dest(mach, &d[TGSI_CHAN_W], &inst->Dst[0], inst, TGSI_CHAN_W, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &d[TGSI_CHAN_W], &inst->Dst[0], inst, TGSI_CHAN_W);
    }
 }
 
@@ -3405,18 +3275,18 @@ exec_log(struct tgsi_exec_machine *mach,
    micro_lg2(&r[1], &r[2]);  /* r1 = lg2(r2) */
    micro_flr(&r[0], &r[1]);  /* r0 = floor(r1) */
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_X) {
-      store_dest(mach, &r[0], &inst->Dst[0], inst, TGSI_CHAN_X, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r[0], &inst->Dst[0], inst, TGSI_CHAN_X);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Y) {
       micro_exp2(&r[0], &r[0]);       /* r0 = 2 ^ r0 */
       micro_div(&r[0], &r[2], &r[0]); /* r0 = r2 / r0 */
-      store_dest(mach, &r[0], &inst->Dst[0], inst, TGSI_CHAN_Y, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r[0], &inst->Dst[0], inst, TGSI_CHAN_Y);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Z) {
-      store_dest(mach, &r[1], &inst->Dst[0], inst, TGSI_CHAN_Z, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r[1], &inst->Dst[0], inst, TGSI_CHAN_Z);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_W) {
-      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_W, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_W);
    }
 }
 
@@ -3430,18 +3300,18 @@ exec_exp(struct tgsi_exec_machine *mach,
    micro_flr(&r[1], &r[0]);  /* r1 = floor(r0) */
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_X) {
       micro_exp2(&r[2], &r[1]);       /* r2 = 2 ^ r1 */
-      store_dest(mach, &r[2], &inst->Dst[0], inst, TGSI_CHAN_X, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r[2], &inst->Dst[0], inst, TGSI_CHAN_X);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Y) {
       micro_sub(&r[2], &r[0], &r[1]); /* r2 = r0 - r1 */
-      store_dest(mach, &r[2], &inst->Dst[0], inst, TGSI_CHAN_Y, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r[2], &inst->Dst[0], inst, TGSI_CHAN_Y);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Z) {
       micro_exp2(&r[2], &r[0]);       /* r2 = 2 ^ r0 */
-      store_dest(mach, &r[2], &inst->Dst[0], inst, TGSI_CHAN_Z, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r[2], &inst->Dst[0], inst, TGSI_CHAN_Z);
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_W) {
-      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_W, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_W);
    }
 }
 
@@ -3463,19 +3333,19 @@ exec_lit(struct tgsi_exec_machine *mach,
          micro_max(&r[2], &r[2], &M128Vec);
          micro_pow(&r[1], &r[1], &r[2]);
          micro_lt(&d[TGSI_CHAN_Z], &ZeroVec, &r[0], &r[1], &ZeroVec);
-         store_dest(mach, &d[TGSI_CHAN_Z], &inst->Dst[0], inst, TGSI_CHAN_Z, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &d[TGSI_CHAN_Z], &inst->Dst[0], inst, TGSI_CHAN_Z);
       }
       if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_Y) {
          micro_max(&d[TGSI_CHAN_Y], &r[0], &ZeroVec);
-         store_dest(mach, &d[TGSI_CHAN_Y], &inst->Dst[0], inst, TGSI_CHAN_Y, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &d[TGSI_CHAN_Y], &inst->Dst[0], inst, TGSI_CHAN_Y);
       }
    }
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_X) {
-      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_X, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_X);
    }
 
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_W) {
-      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_W, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &OneVec, &inst->Dst[0], inst, TGSI_CHAN_W);
    }
 }
 
@@ -3633,9 +3503,9 @@ store_double_channel(struct tgsi_exec_machine *mach,
          }
    }
 
-   store_dest_double(mach, &dst[0], reg, chan_0, TGSI_EXEC_DATA_UINT);
+   store_dest_double(mach, &dst[0], reg, chan_0);
    if (chan_1 != (unsigned)-1)
-      store_dest_double(mach, &dst[1], reg, chan_1, TGSI_EXEC_DATA_UINT);
+      store_dest_double(mach, &dst[1], reg, chan_1);
 }
 
 static void
@@ -3765,7 +3635,7 @@ exec_dfracexp(struct tgsi_exec_machine *mach,
       store_double_channel(mach, &dst, &inst->Dst[0], inst, TGSI_CHAN_Z, TGSI_CHAN_W);
    for (unsigned chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[1].Register.WriteMask & (1 << chan))
-         store_dest(mach, &dst_exp, &inst->Dst[1], inst, chan, TGSI_EXEC_DATA_INT);
+         store_dest(mach, &dst_exp, &inst->Dst[1], inst, chan);
    }
 }
 
@@ -3884,7 +3754,7 @@ exec_load_img(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -3938,7 +3808,7 @@ exec_load_membuf(struct tgsi_exec_machine *mach,
 
    for (int chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &rgba[chan], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &rgba[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -4167,7 +4037,7 @@ exec_atomop_img(struct tgsi_exec_machine *mach,
    }
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -4269,7 +4139,7 @@ exec_atomop_membuf(struct tgsi_exec_machine *mach,
    }
 
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++)
-      store_dest(mach, &r0, &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &r0, &inst->Dst[0], inst, chan);
 }
 
 static void
@@ -4309,8 +4179,7 @@ exec_resq_img(struct tgsi_exec_machine *mach,
 
    for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
       if (inst->Dst[0].Register.WriteMask & (1 << chan)) {
-         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan,
-                    TGSI_EXEC_DATA_INT);
+         store_dest(mach, &r[chan], &inst->Dst[0], inst, chan);
       }
    }
 }
@@ -4329,8 +4198,7 @@ exec_resq_buf(struct tgsi_exec_machine *mach,
 
    if (inst->Dst[0].Register.WriteMask & TGSI_WRITEMASK_X) {
       for (int chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
-         store_dest(mach, &r, &inst->Dst[0], inst, TGSI_CHAN_X,
-                    TGSI_EXEC_DATA_INT);
+         store_dest(mach, &r, &inst->Dst[0], inst, TGSI_CHAN_X);
       }
    }
 }
@@ -4469,8 +4337,7 @@ exec_t_2_64(struct tgsi_exec_machine *mach,
 static void
 exec_64_2_t(struct tgsi_exec_machine *mach,
             const struct tgsi_full_instruction *inst,
-            micro_sop_d op,
-            enum tgsi_exec_datatype dst_datatype)
+            micro_sop_d op)
 {
    union tgsi_double_channel src;
    union tgsi_exec_channel dst;
@@ -4486,7 +4353,7 @@ exec_64_2_t(struct tgsi_exec_machine *mach,
          else
             fetch_double_channel(mach, &src, &inst->Src[0], TGSI_CHAN_Z, TGSI_CHAN_W);
          op(&dst, &src);
-         store_dest(mach, &dst, &inst->Dst[0], inst, bit - 1, dst_datatype);
+         store_dest(mach, &dst, &inst->Dst[0], inst, bit - 1);
       }
    }
 }
@@ -5051,7 +4918,7 @@ exec_interp_at_sample(struct tgsi_exec_machine *mach,
          assert(pos < TGSI_MAX_PRIM_VERTICES * PIPE_MAX_ATTRIBS);
          mach->InputSampleOffsetApply[pos](mach, pos, chan, x, y, &result[chan]);
       }
-      store_dest(mach, &result[chan], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &result[chan], &inst->Dst[0], inst, chan);
    }
 }
 
@@ -5080,7 +4947,7 @@ exec_interp_at_offset(struct tgsi_exec_machine *mach,
       union tgsi_exec_channel result;
       fetch_src_file_channel(mach, TGSI_FILE_INPUT, chan, &index, &index2D, &result);
       mach->InputSampleOffsetApply[pos](mach, pos, chan, ofsx.f[chan], ofsy.f[chan], &result);
-      store_dest(mach, &result, &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &result, &inst->Dst[0], inst, chan);
    }
 }
 
@@ -5116,7 +4983,7 @@ exec_interp_at_centroid(struct tgsi_exec_machine *mach,
        */
       fetch_src_file_channel(mach, TGSI_FILE_INPUT, chan, &index, &index2D,
                              &result[chan]);
-      store_dest(mach, &result[chan], &inst->Dst[0], inst, chan, TGSI_EXEC_DATA_FLOAT);
+      store_dest(mach, &result[chan], &inst->Dst[0], inst, chan);
    }
 
 }
@@ -5139,11 +5006,11 @@ exec_instruction(
 
    switch (inst->Instruction.Opcode) {
    case TGSI_OPCODE_ARL:
-      exec_vector_unary(mach, inst, micro_arl, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_arl, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_MOV:
-      exec_vector_unary(mach, inst, micro_mov, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_mov, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_LIT:
@@ -5151,11 +5018,11 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_RCP:
-      exec_scalar_unary(mach, inst, micro_rcp, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_scalar_unary(mach, inst, micro_rcp, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_RSQ:
-      exec_scalar_unary(mach, inst, micro_rsq, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_scalar_unary(mach, inst, micro_rsq, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_EXP:
@@ -5167,11 +5034,11 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_MUL:
-      exec_vector_binary(mach, inst, micro_mul, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_mul, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_ADD:
-      exec_vector_binary(mach, inst, micro_add, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_add, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_DP3:
@@ -5187,79 +5054,79 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_MIN:
-      exec_vector_binary(mach, inst, micro_min, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_min, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_MAX:
-      exec_vector_binary(mach, inst, micro_max, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_max, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_SLT:
-      exec_vector_binary(mach, inst, micro_slt, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_slt, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_SGE:
-      exec_vector_binary(mach, inst, micro_sge, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_sge, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_MAD:
-      exec_vector_trinary(mach, inst, micro_mad, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_trinary(mach, inst, micro_mad, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_LRP:
-      exec_vector_trinary(mach, inst, micro_lrp, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_trinary(mach, inst, micro_lrp, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_SQRT:
-      exec_scalar_unary(mach, inst, micro_sqrt, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_scalar_unary(mach, inst, micro_sqrt, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_FRC:
-      exec_vector_unary(mach, inst, micro_frc, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_frc, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_FLR:
-      exec_vector_unary(mach, inst, micro_flr, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_flr, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_ROUND:
-      exec_vector_unary(mach, inst, micro_rnd, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_rnd, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_EX2:
-      exec_scalar_unary(mach, inst, micro_exp2, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_scalar_unary(mach, inst, micro_exp2, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_LG2:
-      exec_scalar_unary(mach, inst, micro_lg2, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_scalar_unary(mach, inst, micro_lg2, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_POW:
-      exec_scalar_binary(mach, inst, micro_pow, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_scalar_binary(mach, inst, micro_pow, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_LDEXP:
-      exec_vector_binary(mach, inst, micro_ldexp, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_ldexp, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_COS:
-      exec_scalar_unary(mach, inst, micro_cos, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_scalar_unary(mach, inst, micro_cos, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_DDX_FINE:
-      exec_vector_unary(mach, inst, micro_ddx_fine, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_ddx_fine, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_DDX:
-      exec_vector_unary(mach, inst, micro_ddx, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_ddx, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_DDY_FINE:
-      exec_vector_unary(mach, inst, micro_ddy_fine, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_ddy_fine, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_DDY:
-      exec_vector_unary(mach, inst, micro_ddy, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_ddy, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_KILL:
@@ -5287,23 +5154,23 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_SEQ:
-      exec_vector_binary(mach, inst, micro_seq, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_seq, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_SGT:
-      exec_vector_binary(mach, inst, micro_sgt, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_sgt, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_SIN:
-      exec_scalar_unary(mach, inst, micro_sin, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_scalar_unary(mach, inst, micro_sin, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_SLE:
-      exec_vector_binary(mach, inst, micro_sle, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_sle, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_SNE:
-      exec_vector_binary(mach, inst, micro_sne, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_sne, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_TEX:
@@ -5373,7 +5240,7 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_ARR:
-      exec_vector_unary(mach, inst, micro_arr, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_arr, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_CAL:
@@ -5469,15 +5336,15 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_SSG:
-      exec_vector_unary(mach, inst, micro_sgn, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_sgn, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_CMP:
-      exec_vector_trinary(mach, inst, micro_cmp, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_trinary(mach, inst, micro_cmp, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_DIV:
-      exec_vector_binary(mach, inst, micro_div, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_div, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_DP2:
@@ -5545,39 +5412,39 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_CEIL:
-      exec_vector_unary(mach, inst, micro_ceil, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_ceil, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_I2F:
-      exec_vector_unary(mach, inst, micro_i2f, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_INT);
+      exec_vector_unary(mach, inst, micro_i2f, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_NOT:
-      exec_vector_unary(mach, inst, micro_not, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_unary(mach, inst, micro_not, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_TRUNC:
-      exec_vector_unary(mach, inst, micro_trunc, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_trunc, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_SHL:
-      exec_vector_binary(mach, inst, micro_shl, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_shl, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_AND:
-      exec_vector_binary(mach, inst, micro_and, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_and, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_OR:
-      exec_vector_binary(mach, inst, micro_or, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_or, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_MOD:
-      exec_vector_binary(mach, inst, micro_mod, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_mod, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_XOR:
-      exec_vector_binary(mach, inst, micro_xor, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_xor, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_TXF:
@@ -5685,115 +5552,115 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_F2I:
-      exec_vector_unary(mach, inst, micro_f2i, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_f2i, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_FSEQ:
-      exec_vector_binary(mach, inst, micro_fseq, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_fseq, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_FSGE:
-      exec_vector_binary(mach, inst, micro_fsge, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_fsge, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_FSLT:
-      exec_vector_binary(mach, inst, micro_fslt, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_fslt, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_FSNE:
-      exec_vector_binary(mach, inst, micro_fsne, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_binary(mach, inst, micro_fsne, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_IDIV:
-      exec_vector_binary(mach, inst, micro_idiv, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_idiv, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_IMAX:
-      exec_vector_binary(mach, inst, micro_imax, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_imax, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_IMIN:
-      exec_vector_binary(mach, inst, micro_imin, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_imin, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_INEG:
-      exec_vector_unary(mach, inst, micro_ineg, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_unary(mach, inst, micro_ineg, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_ISGE:
-      exec_vector_binary(mach, inst, micro_isge, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_isge, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_ISHR:
-      exec_vector_binary(mach, inst, micro_ishr, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_ishr, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_ISLT:
-      exec_vector_binary(mach, inst, micro_islt, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_islt, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_F2U:
-      exec_vector_unary(mach, inst, micro_f2u, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_FLOAT);
+      exec_vector_unary(mach, inst, micro_f2u, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_U2F:
-      exec_vector_unary(mach, inst, micro_u2f, TGSI_EXEC_DATA_FLOAT, TGSI_EXEC_DATA_UINT);
+      exec_vector_unary(mach, inst, micro_u2f, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_UADD:
-      exec_vector_binary(mach, inst, micro_uadd, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_uadd, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_UDIV:
-      exec_vector_binary(mach, inst, micro_udiv, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_udiv, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_UMAD:
-      exec_vector_trinary(mach, inst, micro_umad, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_trinary(mach, inst, micro_umad, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_UMAX:
-      exec_vector_binary(mach, inst, micro_umax, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_umax, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_UMIN:
-      exec_vector_binary(mach, inst, micro_umin, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_umin, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_UMOD:
-      exec_vector_binary(mach, inst, micro_umod, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_umod, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_UMUL:
-      exec_vector_binary(mach, inst, micro_umul, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_umul, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_IMUL_HI:
-      exec_vector_binary(mach, inst, micro_imul_hi, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_binary(mach, inst, micro_imul_hi, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_UMUL_HI:
-      exec_vector_binary(mach, inst, micro_umul_hi, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_umul_hi, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_USEQ:
-      exec_vector_binary(mach, inst, micro_useq, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_useq, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_USGE:
-      exec_vector_binary(mach, inst, micro_usge, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_usge, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_USHR:
-      exec_vector_binary(mach, inst, micro_ushr, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_ushr, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_USLT:
-      exec_vector_binary(mach, inst, micro_uslt, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_uslt, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_USNE:
-      exec_vector_binary(mach, inst, micro_usne, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_binary(mach, inst, micro_usne, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_SWITCH:
@@ -5865,7 +5732,7 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_UARL:
-      exec_vector_unary(mach, inst, micro_uarl, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_UINT);
+      exec_vector_unary(mach, inst, micro_uarl, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_UCMP:
@@ -5873,11 +5740,11 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_IABS:
-      exec_vector_unary(mach, inst, micro_iabs, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_unary(mach, inst, micro_iabs, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_ISSG:
-      exec_vector_unary(mach, inst, micro_isgn, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_unary(mach, inst, micro_isgn, TGSI_EXEC_DATA_INT);
       break;
 
    case TGSI_OPCODE_TEX2:
@@ -5903,28 +5770,28 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_IBFE:
-      exec_vector_trinary(mach, inst, micro_ibfe, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_trinary(mach, inst, micro_ibfe, TGSI_EXEC_DATA_INT);
       break;
    case TGSI_OPCODE_UBFE:
-      exec_vector_trinary(mach, inst, micro_ubfe, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_trinary(mach, inst, micro_ubfe, TGSI_EXEC_DATA_UINT);
       break;
    case TGSI_OPCODE_BFI:
-      exec_vector_quaternary(mach, inst, micro_bfi, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_quaternary(mach, inst, micro_bfi, TGSI_EXEC_DATA_UINT);
       break;
    case TGSI_OPCODE_BREV:
-      exec_vector_unary(mach, inst, micro_brev, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_unary(mach, inst, micro_brev, TGSI_EXEC_DATA_UINT);
       break;
    case TGSI_OPCODE_POPC:
-      exec_vector_unary(mach, inst, micro_popc, TGSI_EXEC_DATA_UINT, TGSI_EXEC_DATA_UINT);
+      exec_vector_unary(mach, inst, micro_popc, TGSI_EXEC_DATA_UINT);
       break;
    case TGSI_OPCODE_LSB:
-      exec_vector_unary(mach, inst, micro_lsb, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_UINT);
+      exec_vector_unary(mach, inst, micro_lsb, TGSI_EXEC_DATA_UINT);
       break;
    case TGSI_OPCODE_IMSB:
-      exec_vector_unary(mach, inst, micro_imsb, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_INT);
+      exec_vector_unary(mach, inst, micro_imsb, TGSI_EXEC_DATA_INT);
       break;
    case TGSI_OPCODE_UMSB:
-      exec_vector_unary(mach, inst, micro_umsb, TGSI_EXEC_DATA_INT, TGSI_EXEC_DATA_UINT);
+      exec_vector_unary(mach, inst, micro_umsb, TGSI_EXEC_DATA_UINT);
       break;
 
    case TGSI_OPCODE_F2D:
@@ -5932,7 +5799,7 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_D2F:
-      exec_64_2_t(mach, inst, micro_d2f, TGSI_EXEC_DATA_FLOAT);
+      exec_64_2_t(mach, inst, micro_d2f);
       break;
 
    case TGSI_OPCODE_DABS:
@@ -6012,19 +5879,19 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_I2D:
-      exec_t_2_64(mach, inst, micro_i2d, TGSI_EXEC_DATA_INT);
+      exec_t_2_64(mach, inst, micro_i2d, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_D2I:
-      exec_64_2_t(mach, inst, micro_d2i, TGSI_EXEC_DATA_INT);
+      exec_64_2_t(mach, inst, micro_d2i);
       break;
 
    case TGSI_OPCODE_U2D:
-      exec_t_2_64(mach, inst, micro_u2d, TGSI_EXEC_DATA_UINT);
+      exec_t_2_64(mach, inst, micro_u2d, TGSI_EXEC_DATA_FLOAT);
       break;
 
    case TGSI_OPCODE_D2U:
-      exec_64_2_t(mach, inst, micro_d2u, TGSI_EXEC_DATA_INT);
+      exec_64_2_t(mach, inst, micro_d2u);
       break;
 
    case TGSI_OPCODE_LOAD:
@@ -6155,10 +6022,10 @@ exec_instruction(
       break;
 
    case TGSI_OPCODE_U642F:
-      exec_64_2_t(mach, inst, micro_u642f, TGSI_EXEC_DATA_FLOAT);
+      exec_64_2_t(mach, inst, micro_u642f);
       break;
    case TGSI_OPCODE_I642F:
-      exec_64_2_t(mach, inst, micro_i642f, TGSI_EXEC_DATA_FLOAT);
+      exec_64_2_t(mach, inst, micro_i642f);
       break;
 
    case TGSI_OPCODE_U642D:
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_exec.h b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_exec.h
index 982cad592..257769819 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_exec.h
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_exec.h
@@ -76,12 +76,12 @@ union tgsi_exec_channel
    float    f[TGSI_QUAD_SIZE];
    int      i[TGSI_QUAD_SIZE];
    unsigned u[TGSI_QUAD_SIZE];
-};
+} ALIGN16;
 
 /**
   * A vector[RGBA] of channels[4 pixels]
   */
-struct tgsi_exec_vector
+struct ALIGN16 tgsi_exec_vector
 {
    union tgsi_exec_channel xyzw[TGSI_NUM_CHANNELS];
 };
@@ -286,7 +286,7 @@ typedef void (* apply_sample_offset_func)(
 /**
  * Run-time virtual machine state for executing TGSI shader.
  */
-struct tgsi_exec_machine
+struct ALIGN16 tgsi_exec_machine
 {
    /* Total = program temporaries + internal temporaries
     */
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_from_mesa.c b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_from_mesa.c
index 4b4882765..9e3949711 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_from_mesa.c
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_from_mesa.c
@@ -265,11 +265,11 @@ tgsi_get_sysval_semantic(unsigned sysval)
    /* Compute shader */
    case SYSTEM_VALUE_LOCAL_INVOCATION_ID:
       return TGSI_SEMANTIC_THREAD_ID;
-   case SYSTEM_VALUE_WORK_GROUP_ID:
+   case SYSTEM_VALUE_WORKGROUP_ID:
       return TGSI_SEMANTIC_BLOCK_ID;
-   case SYSTEM_VALUE_NUM_WORK_GROUPS:
+   case SYSTEM_VALUE_NUM_WORKGROUPS:
       return TGSI_SEMANTIC_GRID_SIZE;
-   case SYSTEM_VALUE_LOCAL_GROUP_SIZE:
+   case SYSTEM_VALUE_WORKGROUP_SIZE:
       return TGSI_SEMANTIC_BLOCK_SIZE;
 
    /* ARB_shader_ballot */
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_lowering.c b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_lowering.c
index 4f8f30175..db2302669 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_lowering.c
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_lowering.c
@@ -1188,7 +1188,6 @@ emit_twoside(struct tgsi_transform_context *tctx)
       decl.Declaration.Interpolate = true;
       decl.Interp.Interpolate = info->input_interpolate[in_idx];
       decl.Interp.Location = info->input_interpolate_loc[in_idx];
-      decl.Interp.CylindricalWrap = info->input_cylindrical_wrap[in_idx];
       tctx->emit_declaration(tctx, &decl);
    }
 
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_scan.c b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_scan.c
index 0cd4a8883..ecb3706ed 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_scan.c
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_scan.c
@@ -668,7 +668,6 @@ scan_declaration(struct tgsi_shader_info *info,
          info->input_semantic_index[reg] = (ubyte) semIndex;
          info->input_interpolate[reg] = (ubyte)fulldecl->Interp.Interpolate;
          info->input_interpolate_loc[reg] = (ubyte)fulldecl->Interp.Location;
-         info->input_cylindrical_wrap[reg] = (ubyte)fulldecl->Interp.CylindricalWrap;
 
          /* Vertex shaders can have inputs with holes between them. */
          info->num_inputs = MAX2(info->num_inputs, reg + 1);
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_scan.h b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_scan.h
index ca8d90a80..ace5b0887 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_scan.h
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_scan.h
@@ -51,7 +51,6 @@ struct tgsi_shader_info
    ubyte input_interpolate[PIPE_MAX_SHADER_INPUTS];
    ubyte input_interpolate_loc[PIPE_MAX_SHADER_INPUTS];
    ubyte input_usage_mask[PIPE_MAX_SHADER_INPUTS];
-   ubyte input_cylindrical_wrap[PIPE_MAX_SHADER_INPUTS];
    ubyte output_semantic_name[PIPE_MAX_SHADER_OUTPUTS]; /**< TGSI_SEMANTIC_x */
    ubyte output_semantic_index[PIPE_MAX_SHADER_OUTPUTS];
    ubyte output_usagemask[PIPE_MAX_SHADER_OUTPUTS];
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_ureg.c b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_ureg.c
index 77444a1ca..eedba34cb 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_ureg.c
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_ureg.c
@@ -114,11 +114,10 @@ struct ureg_program
    bool supports_any_inout_decl_range;
    int next_shader_processor;
 
-   struct {
+   struct ureg_input_decl {
       enum tgsi_semantic semantic_name;
       unsigned semantic_index;
       enum tgsi_interpolate_mode interp;
-      unsigned char cylindrical_wrap;
       unsigned char usage_mask;
       enum tgsi_interpolate_loc interp_location;
       unsigned first;
@@ -135,7 +134,7 @@ struct ureg_program
    } system_value[UREG_MAX_SYSTEM_VALUE];
    unsigned nr_system_values;
 
-   struct {
+   struct ureg_output_decl {
       enum tgsi_semantic semantic_name;
       unsigned semantic_index;
       unsigned streams;
@@ -284,11 +283,10 @@ ureg_property(struct ureg_program *ureg, unsigned name, unsigned value)
 }
 
 struct ureg_src
-ureg_DECL_fs_input_cyl_centroid_layout(struct ureg_program *ureg,
+ureg_DECL_fs_input_centroid_layout(struct ureg_program *ureg,
                        enum tgsi_semantic semantic_name,
                        unsigned semantic_index,
                        enum tgsi_interpolate_mode interp_mode,
-                       unsigned cylindrical_wrap,
                        enum tgsi_interpolate_loc interp_location,
                        unsigned index,
                        unsigned usage_mask,
@@ -304,7 +302,6 @@ ureg_DECL_fs_input_cyl_centroid_layout(struct ureg_program *ureg,
       if (ureg->input[i].semantic_name == semantic_name &&
           ureg->input[i].semantic_index == semantic_index) {
          assert(ureg->input[i].interp == interp_mode);
-         assert(ureg->input[i].cylindrical_wrap == cylindrical_wrap);
          assert(ureg->input[i].interp_location == interp_location);
          if (ureg->input[i].array_id == array_id) {
             ureg->input[i].usage_mask |= usage_mask;
@@ -319,7 +316,6 @@ ureg_DECL_fs_input_cyl_centroid_layout(struct ureg_program *ureg,
       ureg->input[i].semantic_name = semantic_name;
       ureg->input[i].semantic_index = semantic_index;
       ureg->input[i].interp = interp_mode;
-      ureg->input[i].cylindrical_wrap = cylindrical_wrap;
       ureg->input[i].interp_location = interp_location;
       ureg->input[i].first = index;
       ureg->input[i].last = index + array_size - 1;
@@ -337,18 +333,17 @@ out:
 }
 
 struct ureg_src
-ureg_DECL_fs_input_cyl_centroid(struct ureg_program *ureg,
+ureg_DECL_fs_input_centroid(struct ureg_program *ureg,
                        enum tgsi_semantic semantic_name,
                        unsigned semantic_index,
                        enum tgsi_interpolate_mode interp_mode,
-                       unsigned cylindrical_wrap,
                        enum tgsi_interpolate_loc interp_location,
                        unsigned array_id,
                        unsigned array_size)
 {
-   return ureg_DECL_fs_input_cyl_centroid_layout(ureg,
+   return ureg_DECL_fs_input_centroid_layout(ureg,
          semantic_name, semantic_index, interp_mode,
-         cylindrical_wrap, interp_location,
+         interp_location,
          ureg->nr_input_regs, TGSI_WRITEMASK_XYZW, array_id, array_size);
 }
 
@@ -374,9 +369,9 @@ ureg_DECL_input_layout(struct ureg_program *ureg,
                 unsigned array_id,
                 unsigned array_size)
 {
-   return ureg_DECL_fs_input_cyl_centroid_layout(ureg,
+   return ureg_DECL_fs_input_centroid_layout(ureg,
                semantic_name, semantic_index,
-               TGSI_INTERPOLATE_CONSTANT, 0, TGSI_INTERPOLATE_LOC_CENTER,
+               TGSI_INTERPOLATE_CONSTANT, TGSI_INTERPOLATE_LOC_CENTER,
                index, usage_mask, array_id, array_size);
 }
 
@@ -388,8 +383,8 @@ ureg_DECL_input(struct ureg_program *ureg,
                 unsigned array_id,
                 unsigned array_size)
 {
-   return ureg_DECL_fs_input_cyl_centroid(ureg, semantic_name, semantic_index,
-                                          TGSI_INTERPOLATE_CONSTANT, 0,
+   return ureg_DECL_fs_input_centroid(ureg, semantic_name, semantic_index,
+                                          TGSI_INTERPOLATE_CONSTANT,
                                           TGSI_INTERPOLATE_LOC_CENTER,
                                           array_id, array_size);
 }
@@ -1587,7 +1582,6 @@ emit_decl_fs(struct ureg_program *ureg,
              enum tgsi_semantic semantic_name,
              unsigned semantic_index,
              enum tgsi_interpolate_mode interpolate,
-             unsigned cylindrical_wrap,
              enum tgsi_interpolate_loc interpolate_location,
              unsigned array_id,
              unsigned usage_mask)
@@ -1610,7 +1604,6 @@ emit_decl_fs(struct ureg_program *ureg,
 
    out[2].value = 0;
    out[2].decl_interp.Interpolate = interpolate;
-   out[2].decl_interp.CylindricalWrap = cylindrical_wrap;
    out[2].decl_interp.Location = interpolate_location;
 
    out[3].value = 0;
@@ -1819,6 +1812,21 @@ emit_property(struct ureg_program *ureg,
    out[1].prop_data.Data = data;
 }
 
+static int
+input_sort(const void *in_a, const void *in_b)
+{
+   const struct ureg_input_decl *a = in_a, *b = in_b;
+
+   return a->first - b->first;
+}
+
+static int
+output_sort(const void *in_a, const void *in_b)
+{
+   const struct ureg_output_decl *a = in_a, *b = in_b;
+
+   return a->first - b->first;
+}
 
 static void emit_decls( struct ureg_program *ureg )
 {
@@ -1828,6 +1836,11 @@ static void emit_decls( struct ureg_program *ureg )
       if (ureg->properties[i] != ~0u)
          emit_property(ureg, i, ureg->properties[i]);
 
+   /* While not required by TGSI spec, virglrenderer has a dependency on the
+    * inputs being sorted.
+    */
+   qsort(ureg->input, ureg->nr_inputs, sizeof(ureg->input[0]), input_sort);
+
    if (ureg->processor == PIPE_SHADER_VERTEX) {
       for (i = 0; i < PIPE_MAX_ATTRIBS; i++) {
          if (ureg->vs_inputs[i/32] & (1u << (i%32))) {
@@ -1844,7 +1857,6 @@ static void emit_decls( struct ureg_program *ureg )
                          ureg->input[i].semantic_name,
                          ureg->input[i].semantic_index,
                          ureg->input[i].interp,
-                         ureg->input[i].cylindrical_wrap,
                          ureg->input[i].interp_location,
                          ureg->input[i].array_id,
                          ureg->input[i].usage_mask);
@@ -1860,7 +1872,6 @@ static void emit_decls( struct ureg_program *ureg )
                             ureg->input[i].semantic_index +
                             (j - ureg->input[i].first),
                             ureg->input[i].interp,
-                            ureg->input[i].cylindrical_wrap,
                             ureg->input[i].interp_location, 0,
                             ureg->input[i].usage_mask);
             }
@@ -1908,6 +1919,11 @@ static void emit_decls( struct ureg_program *ureg )
                          TGSI_WRITEMASK_XYZW, 0, FALSE);
    }
 
+   /* While not required by TGSI spec, virglrenderer has a dependency on the
+    * outputs being sorted.
+    */
+   qsort(ureg->output, ureg->nr_outputs, sizeof(ureg->output[0]), output_sort);
+
    if (ureg->supports_any_inout_decl_range) {
       for (i = 0; i < ureg->nr_outputs; i++) {
          emit_decl_semantic(ureg,
@@ -2347,6 +2363,11 @@ ureg_setup_fragment_shader(struct ureg_program *ureg,
          assert(0);
       }
    }
+
+   if (info->fs.advanced_blend_modes) {
+      ureg_property(ureg, TGSI_PROPERTY_FS_BLEND_EQUATION_ADVANCED,
+                    info->fs.advanced_blend_modes);
+   }
 }
 
 static void
@@ -2354,11 +2375,11 @@ ureg_setup_compute_shader(struct ureg_program *ureg,
                           const struct shader_info *info)
 {
    ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH,
-                 info->cs.local_size[0]);
+                 info->workgroup_size[0]);
    ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT,
-                 info->cs.local_size[1]);
+                 info->workgroup_size[1]);
    ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH,
-                 info->cs.local_size[2]);
+                 info->workgroup_size[2]);
 
    if (info->shared_size)
       ureg_DECL_memory(ureg, TGSI_MEMORY_TYPE_SHARED);
diff --git a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_ureg.h b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_ureg.h
index 1b69c8730..343708b6c 100644
--- a/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_ureg.h
+++ b/lib/mesa/src/gallium/auxiliary/tgsi/tgsi_ureg.h
@@ -172,11 +172,10 @@ ureg_property(struct ureg_program *ureg, unsigned name, unsigned value);
  */
 
 struct ureg_src
-ureg_DECL_fs_input_cyl_centroid_layout(struct ureg_program *,
+ureg_DECL_fs_input_centroid_layout(struct ureg_program *,
                        enum tgsi_semantic semantic_name,
                        unsigned semantic_index,
                        enum tgsi_interpolate_mode interp_mode,
-                       unsigned cylindrical_wrap,
                        enum tgsi_interpolate_loc interp_location,
                        unsigned index,
                        unsigned usage_mask,
@@ -184,41 +183,25 @@ ureg_DECL_fs_input_cyl_centroid_layout(struct ureg_program *,
                        unsigned array_size);
 
 struct ureg_src
-ureg_DECL_fs_input_cyl_centroid(struct ureg_program *,
+ureg_DECL_fs_input_centroid(struct ureg_program *,
                        enum tgsi_semantic semantic_name,
                        unsigned semantic_index,
                        enum tgsi_interpolate_mode interp_mode,
-                       unsigned cylindrical_wrap,
                        enum tgsi_interpolate_loc interp_location,
                        unsigned array_id,
                        unsigned array_size);
 
 static inline struct ureg_src
-ureg_DECL_fs_input_cyl(struct ureg_program *ureg,
-                       enum tgsi_semantic semantic_name,
-                       unsigned semantic_index,
-                       enum tgsi_interpolate_mode interp_mode,
-                       unsigned cylindrical_wrap)
-{
-   return ureg_DECL_fs_input_cyl_centroid(ureg,
-                                 semantic_name,
-                                 semantic_index,
-                                 interp_mode,
-                                 cylindrical_wrap,
-                                 TGSI_INTERPOLATE_LOC_CENTER, 0, 1);
-}
-
-static inline struct ureg_src
 ureg_DECL_fs_input(struct ureg_program *ureg,
                    enum tgsi_semantic semantic_name,
                    unsigned semantic_index,
                    enum tgsi_interpolate_mode interp_mode)
 {
-   return ureg_DECL_fs_input_cyl_centroid(ureg,
+   return ureg_DECL_fs_input_centroid(ureg,
                                  semantic_name,
                                  semantic_index,
                                  interp_mode,
-                                 0, TGSI_INTERPOLATE_LOC_CENTER, 0, 1);
+                                 TGSI_INTERPOLATE_LOC_CENTER, 0, 1);
 }
 
 struct ureg_src