radeonsi: have separate LS and ES main shader parts in the shader selector
[mesa.git] / src / gallium / drivers / radeonsi / si_shader.c
index daff6296c948b742fe621a49bb861986425b6c56..de427789aae8aa5dc3df534d112e5cd75040714a 100644 (file)
@@ -98,14 +98,6 @@ enum {
        LOCAL_ADDR_SPACE = 3,
 };
 
-#define SENDMSG_GS 2
-#define SENDMSG_GS_DONE 3
-
-#define SENDMSG_GS_OP_NOP      (0 << 4)
-#define SENDMSG_GS_OP_CUT      (1 << 4)
-#define SENDMSG_GS_OP_EMIT     (2 << 4)
-#define SENDMSG_GS_OP_EMIT_CUT (3 << 4)
-
 /**
  * Returns a unique index for a semantic name and index. The index must be
  * less than 64, so that a 64-bit bitmask of used inputs or outputs can be
@@ -176,7 +168,7 @@ static LLVMValueRef unpack_param(struct si_shader_context *ctx,
                                          param);
 
        if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMFloatTypeKind)
-               value = bitcast(&ctx->soa.bld_base,
+               value = bitcast(&ctx->bld_base,
                                TGSI_TYPE_UNSIGNED, value);
 
        if (rshift)
@@ -251,7 +243,7 @@ get_tcs_out_patch_stride(struct si_shader_context *ctx)
 static LLVMValueRef
 get_tcs_out_patch0_offset(struct si_shader_context *ctx)
 {
-       return lp_build_mul_imm(&ctx->soa.bld_base.uint_bld,
+       return lp_build_mul_imm(&ctx->bld_base.uint_bld,
                                unpack_param(ctx,
                                             SI_PARAM_TCS_OUT_OFFSETS,
                                             0, 16),
@@ -261,7 +253,7 @@ get_tcs_out_patch0_offset(struct si_shader_context *ctx)
 static LLVMValueRef
 get_tcs_out_patch0_patch_data_offset(struct si_shader_context *ctx)
 {
-       return lp_build_mul_imm(&ctx->soa.bld_base.uint_bld,
+       return lp_build_mul_imm(&ctx->bld_base.uint_bld,
                                unpack_param(ctx,
                                             SI_PARAM_TCS_OUT_OFFSETS,
                                             16, 16),
@@ -307,71 +299,13 @@ get_tcs_out_current_patch_data_offset(struct si_shader_context *ctx)
                            "");
 }
 
-static LLVMValueRef build_gep0(struct si_shader_context *ctx,
-                              LLVMValueRef base_ptr, LLVMValueRef index)
-{
-       LLVMValueRef indices[2] = {
-               LLVMConstInt(ctx->i32, 0, 0),
-               index,
-       };
-       return LLVMBuildGEP(ctx->gallivm.builder, base_ptr,
-                           indices, 2, "");
-}
-
-static void build_indexed_store(struct si_shader_context *ctx,
-                               LLVMValueRef base_ptr, LLVMValueRef index,
-                               LLVMValueRef value)
-{
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-
-       LLVMBuildStore(gallivm->builder, value,
-                      build_gep0(ctx, base_ptr, index));
-}
-
-/**
- * Build an LLVM bytecode indexed load using LLVMBuildGEP + LLVMBuildLoad.
- * It's equivalent to doing a load from &base_ptr[index].
- *
- * \param base_ptr  Where the array starts.
- * \param index     The element index into the array.
- * \param uniform   Whether the base_ptr and index can be assumed to be
- *                  dynamically uniform
- */
-static LLVMValueRef build_indexed_load(struct si_shader_context *ctx,
-                                      LLVMValueRef base_ptr, LLVMValueRef index,
-                                      bool uniform)
-{
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef pointer;
-
-       pointer = build_gep0(ctx, base_ptr, index);
-       if (uniform)
-               LLVMSetMetadata(pointer, ctx->uniform_md_kind, ctx->empty_md);
-       return LLVMBuildLoad(gallivm->builder, pointer, "");
-}
-
-/**
- * Do a load from &base_ptr[index], but also add a flag that it's loading
- * a constant from a dynamically uniform index.
- */
-static LLVMValueRef build_indexed_load_const(
-       struct si_shader_context *ctx,
-       LLVMValueRef base_ptr, LLVMValueRef index)
-{
-       LLVMValueRef result = build_indexed_load(ctx, base_ptr, index, true);
-       LLVMSetMetadata(result, ctx->invariant_load_md_kind, ctx->empty_md);
-       return result;
-}
-
 static LLVMValueRef get_instance_index_for_fetch(
        struct si_shader_context *radeon_bld,
        unsigned param_start_instance, unsigned divisor)
 {
        struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->soa.bld_base);
-       struct gallivm_state *gallivm = radeon_bld->soa.bld_base.base.gallivm;
+               si_shader_context(&radeon_bld->bld_base);
+       struct gallivm_state *gallivm = radeon_bld->bld_base.base.gallivm;
 
        LLVMValueRef result = LLVMGetParam(radeon_bld->main_fn,
                                           ctx->param_instance_id);
@@ -385,56 +319,102 @@ static LLVMValueRef get_instance_index_for_fetch(
                            LLVMGetParam(radeon_bld->main_fn, param_start_instance), "");
 }
 
+/* Bitcast <4 x float> to <2 x double>, extract the component, and convert
+ * to float. */
+static LLVMValueRef extract_double_to_float(struct si_shader_context *ctx,
+                                           LLVMValueRef vec4,
+                                           unsigned double_index)
+{
+       LLVMBuilderRef builder = ctx->gallivm.builder;
+       LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->gallivm.context);
+       LLVMValueRef dvec2 = LLVMBuildBitCast(builder, vec4,
+                                             LLVMVectorType(f64, 2), "");
+       LLVMValueRef index = LLVMConstInt(ctx->i32, double_index, 0);
+       LLVMValueRef value = LLVMBuildExtractElement(builder, dvec2, index, "");
+       return LLVMBuildFPTrunc(builder, value, ctx->f32, "");
+}
+
 static void declare_input_vs(
        struct si_shader_context *ctx,
        unsigned input_index,
        const struct tgsi_full_declaration *decl,
        LLVMValueRef out[4])
 {
-       struct lp_build_context *base = &ctx->soa.bld_base.base;
+       struct lp_build_context *base = &ctx->bld_base.base;
        struct gallivm_state *gallivm = base->gallivm;
 
        unsigned chan;
        unsigned fix_fetch;
+       unsigned num_fetches;
+       unsigned fetch_stride;
 
        LLVMValueRef t_list_ptr;
        LLVMValueRef t_offset;
        LLVMValueRef t_list;
-       LLVMValueRef attribute_offset;
-       LLVMValueRef buffer_index;
+       LLVMValueRef vertex_index;
        LLVMValueRef args[3];
-       LLVMValueRef input;
+       LLVMValueRef input[3];
 
        /* Load the T list */
        t_list_ptr = LLVMGetParam(ctx->main_fn, SI_PARAM_VERTEX_BUFFERS);
 
        t_offset = lp_build_const_int32(gallivm, input_index);
 
-       t_list = build_indexed_load_const(ctx, t_list_ptr, t_offset);
+       t_list = ac_build_indexed_load_const(&ctx->ac, t_list_ptr, t_offset);
 
-       /* Build the attribute offset */
-       attribute_offset = lp_build_const_int32(gallivm, 0);
-
-       buffer_index = LLVMGetParam(ctx->main_fn,
+       vertex_index = LLVMGetParam(ctx->main_fn,
                                    ctx->param_vertex_index0 +
                                    input_index);
 
+       fix_fetch = ctx->shader->key.mono.vs.fix_fetch[input_index];
+
+       /* Do multiple loads for special formats. */
+       switch (fix_fetch) {
+       case SI_FIX_FETCH_RGB_64_FLOAT:
+               num_fetches = 3; /* 3 2-dword loads */
+               fetch_stride = 8;
+               break;
+       case SI_FIX_FETCH_RGBA_64_FLOAT:
+               num_fetches = 2; /* 2 4-dword loads */
+               fetch_stride = 16;
+               break;
+       case SI_FIX_FETCH_RGB_8:
+       case SI_FIX_FETCH_RGB_8_INT:
+               num_fetches = 3;
+               fetch_stride = 1;
+               break;
+       case SI_FIX_FETCH_RGB_16:
+       case SI_FIX_FETCH_RGB_16_INT:
+               num_fetches = 3;
+               fetch_stride = 2;
+               break;
+       default:
+               num_fetches = 1;
+               fetch_stride = 0;
+       }
+
        args[0] = t_list;
-       args[1] = attribute_offset;
-       args[2] = buffer_index;
-       input = lp_build_intrinsic(gallivm->builder,
-               "llvm.SI.vs.load.input", ctx->v4f32, args, 3,
-               LP_FUNC_ATTR_READNONE);
+       args[2] = vertex_index;
+
+       for (unsigned i = 0; i < num_fetches; i++) {
+               args[1] = LLVMConstInt(ctx->i32, fetch_stride * i, 0);
+
+               input[i] = lp_build_intrinsic(gallivm->builder,
+                       "llvm.SI.vs.load.input", ctx->v4f32, args, 3,
+                       LP_FUNC_ATTR_READNONE);
+       }
 
        /* Break up the vec4 into individual components */
        for (chan = 0; chan < 4; chan++) {
                LLVMValueRef llvm_chan = lp_build_const_int32(gallivm, chan);
                out[chan] = LLVMBuildExtractElement(gallivm->builder,
-                                                   input, llvm_chan, "");
+                                                   input[0], llvm_chan, "");
        }
 
-       fix_fetch = (ctx->shader->key.mono.vs.fix_fetch >> (2 * input_index)) & 3;
-       if (fix_fetch) {
+       switch (fix_fetch) {
+       case SI_FIX_FETCH_A2_SNORM:
+       case SI_FIX_FETCH_A2_SSCALED:
+       case SI_FIX_FETCH_A2_SINT: {
                /* The hardware returns an unsigned value; convert it to a
                 * signed one.
                 */
@@ -470,6 +450,98 @@ static void declare_input_vs(
                }
 
                out[3] = tmp;
+               break;
+       }
+       case SI_FIX_FETCH_RGBA_32_UNORM:
+       case SI_FIX_FETCH_RGBX_32_UNORM:
+               for (chan = 0; chan < 4; chan++) {
+                       out[chan] = LLVMBuildBitCast(gallivm->builder, out[chan],
+                                                    ctx->i32, "");
+                       out[chan] = LLVMBuildUIToFP(gallivm->builder,
+                                                   out[chan], ctx->f32, "");
+                       out[chan] = LLVMBuildFMul(gallivm->builder, out[chan],
+                                                 LLVMConstReal(ctx->f32, 1.0 / UINT_MAX), "");
+               }
+               /* RGBX UINT returns 1 in alpha, which would be rounded to 0 by normalizing. */
+               if (fix_fetch == SI_FIX_FETCH_RGBX_32_UNORM)
+                       out[3] = LLVMConstReal(ctx->f32, 1);
+               break;
+       case SI_FIX_FETCH_RGBA_32_SNORM:
+       case SI_FIX_FETCH_RGBX_32_SNORM:
+       case SI_FIX_FETCH_RGBA_32_FIXED:
+       case SI_FIX_FETCH_RGBX_32_FIXED: {
+               double scale;
+               if (fix_fetch >= SI_FIX_FETCH_RGBA_32_FIXED)
+                       scale = 1.0 / 0x10000;
+               else
+                       scale = 1.0 / INT_MAX;
+
+               for (chan = 0; chan < 4; chan++) {
+                       out[chan] = LLVMBuildBitCast(gallivm->builder, out[chan],
+                                                    ctx->i32, "");
+                       out[chan] = LLVMBuildSIToFP(gallivm->builder,
+                                                   out[chan], ctx->f32, "");
+                       out[chan] = LLVMBuildFMul(gallivm->builder, out[chan],
+                                                 LLVMConstReal(ctx->f32, scale), "");
+               }
+               /* RGBX SINT returns 1 in alpha, which would be rounded to 0 by normalizing. */
+               if (fix_fetch == SI_FIX_FETCH_RGBX_32_SNORM ||
+                   fix_fetch == SI_FIX_FETCH_RGBX_32_FIXED)
+                       out[3] = LLVMConstReal(ctx->f32, 1);
+               break;
+       }
+       case SI_FIX_FETCH_RGBA_32_USCALED:
+               for (chan = 0; chan < 4; chan++) {
+                       out[chan] = LLVMBuildBitCast(gallivm->builder, out[chan],
+                                                    ctx->i32, "");
+                       out[chan] = LLVMBuildUIToFP(gallivm->builder,
+                                                   out[chan], ctx->f32, "");
+               }
+               break;
+       case SI_FIX_FETCH_RGBA_32_SSCALED:
+               for (chan = 0; chan < 4; chan++) {
+                       out[chan] = LLVMBuildBitCast(gallivm->builder, out[chan],
+                                                    ctx->i32, "");
+                       out[chan] = LLVMBuildSIToFP(gallivm->builder,
+                                                   out[chan], ctx->f32, "");
+               }
+               break;
+       case SI_FIX_FETCH_RG_64_FLOAT:
+               for (chan = 0; chan < 2; chan++)
+                       out[chan] = extract_double_to_float(ctx, input[0], chan);
+
+               out[2] = LLVMConstReal(ctx->f32, 0);
+               out[3] = LLVMConstReal(ctx->f32, 1);
+               break;
+       case SI_FIX_FETCH_RGB_64_FLOAT:
+               for (chan = 0; chan < 3; chan++)
+                       out[chan] = extract_double_to_float(ctx, input[chan], 0);
+
+               out[3] = LLVMConstReal(ctx->f32, 1);
+               break;
+       case SI_FIX_FETCH_RGBA_64_FLOAT:
+               for (chan = 0; chan < 4; chan++) {
+                       out[chan] = extract_double_to_float(ctx, input[chan / 2],
+                                                           chan % 2);
+               }
+               break;
+       case SI_FIX_FETCH_RGB_8:
+       case SI_FIX_FETCH_RGB_8_INT:
+       case SI_FIX_FETCH_RGB_16:
+       case SI_FIX_FETCH_RGB_16_INT:
+               for (chan = 0; chan < 3; chan++) {
+                       out[chan] = LLVMBuildExtractElement(gallivm->builder,
+                                                           input[chan],
+                                                           ctx->i32_0, "");
+               }
+               if (fix_fetch == SI_FIX_FETCH_RGB_8 ||
+                   fix_fetch == SI_FIX_FETCH_RGB_16) {
+                       out[3] = LLVMConstReal(ctx->f32, 1);
+               } else {
+                       out[3] = LLVMBuildBitCast(gallivm->builder, ctx->i32_1,
+                                                 ctx->f32, "");
+               }
+               break;
        }
 }
 
@@ -508,10 +580,10 @@ static LLVMValueRef get_indirect_index(struct si_shader_context *ctx,
                                       const struct tgsi_ind_register *ind,
                                       int rel_index)
 {
-       struct gallivm_state *gallivm = ctx->soa.bld_base.base.gallivm;
+       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
        LLVMValueRef result;
 
-       result = ctx->soa.addr[ind->Index][ind->Swizzle];
+       result = ctx->addrs[ind->Index][ind->Swizzle];
        result = LLVMBuildLoad(gallivm->builder, result, "");
        result = LLVMBuildAdd(gallivm->builder, result,
                              lp_build_const_int32(gallivm, rel_index), "");
@@ -548,7 +620,7 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
                                   LLVMValueRef vertex_dw_stride,
                                   LLVMValueRef base_addr)
 {
-       struct gallivm_state *gallivm = ctx->soa.bld_base.base.gallivm;
+       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
        struct tgsi_shader_info *info = &ctx->shader->selector->info;
        ubyte *name, *index, *array_first;
        int first, param;
@@ -646,7 +718,7 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx,
                                                LLVMValueRef vertex_index,
                                                LLVMValueRef param_index)
 {
-       struct gallivm_state *gallivm = ctx->soa.bld_base.base.gallivm;
+       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
        LLVMValueRef base_addr, vertices_per_patch, num_patches, total_vertices;
        LLVMValueRef param_stride, constant16;
 
@@ -690,7 +762,7 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
                                        const struct tgsi_full_dst_register *dst,
                                        const struct tgsi_full_src_register *src)
 {
-       struct gallivm_state *gallivm = ctx->soa.bld_base.base.gallivm;
+       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
        struct tgsi_shader_info *info = &ctx->shader->selector->info;
        ubyte *name, *index, *array_first;
        struct tgsi_full_src_register reg;
@@ -748,151 +820,6 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
        return get_tcs_tes_buffer_address(ctx, vertex_index, param_index);
 }
 
-/* TBUFFER_STORE_FORMAT_{X,XY,XYZ,XYZW} <- the suffix is selected by num_channels=1..4.
- * The type of vdata must be one of i32 (num_channels=1), v2i32 (num_channels=2),
- * or v4i32 (num_channels=3,4). */
-static void build_tbuffer_store(struct si_shader_context *ctx,
-                               LLVMValueRef rsrc,
-                               LLVMValueRef vdata,
-                               unsigned num_channels,
-                               LLVMValueRef vaddr,
-                               LLVMValueRef soffset,
-                               unsigned inst_offset,
-                               unsigned dfmt,
-                               unsigned nfmt,
-                               unsigned offen,
-                               unsigned idxen,
-                               unsigned glc,
-                               unsigned slc,
-                               unsigned tfe)
-{
-       struct gallivm_state *gallivm = &ctx->gallivm;
-       LLVMValueRef args[] = {
-               rsrc,
-               vdata,
-               LLVMConstInt(ctx->i32, num_channels, 0),
-               vaddr,
-               soffset,
-               LLVMConstInt(ctx->i32, inst_offset, 0),
-               LLVMConstInt(ctx->i32, dfmt, 0),
-               LLVMConstInt(ctx->i32, nfmt, 0),
-               LLVMConstInt(ctx->i32, offen, 0),
-               LLVMConstInt(ctx->i32, idxen, 0),
-               LLVMConstInt(ctx->i32, glc, 0),
-               LLVMConstInt(ctx->i32, slc, 0),
-               LLVMConstInt(ctx->i32, tfe, 0)
-       };
-
-       /* The instruction offset field has 12 bits */
-       assert(offen || inst_offset < (1 << 12));
-
-       /* The intrinsic is overloaded, we need to add a type suffix for overloading to work. */
-       unsigned func = CLAMP(num_channels, 1, 3) - 1;
-       const char *types[] = {"i32", "v2i32", "v4i32"};
-       char name[256];
-       snprintf(name, sizeof(name), "llvm.SI.tbuffer.store.%s", types[func]);
-
-       lp_build_intrinsic(gallivm->builder, name, ctx->voidt,
-                          args, ARRAY_SIZE(args), 0);
-}
-
-static void build_tbuffer_store_dwords(struct si_shader_context *ctx,
-                                    LLVMValueRef rsrc,
-                                    LLVMValueRef vdata,
-                                    unsigned num_channels,
-                                    LLVMValueRef vaddr,
-                                    LLVMValueRef soffset,
-                                    unsigned inst_offset)
-{
-       static unsigned dfmt[] = {
-               V_008F0C_BUF_DATA_FORMAT_32,
-               V_008F0C_BUF_DATA_FORMAT_32_32,
-               V_008F0C_BUF_DATA_FORMAT_32_32_32,
-               V_008F0C_BUF_DATA_FORMAT_32_32_32_32
-       };
-       assert(num_channels >= 1 && num_channels <= 4);
-
-       build_tbuffer_store(ctx, rsrc, vdata, num_channels, vaddr, soffset,
-                           inst_offset, dfmt[num_channels-1],
-                           V_008F0C_BUF_NUM_FORMAT_UINT, 1, 0, 1, 1, 0);
-}
-
-static LLVMValueRef build_buffer_load(struct si_shader_context *ctx,
-                                      LLVMValueRef rsrc,
-                                      int num_channels,
-                                      LLVMValueRef vindex,
-                                      LLVMValueRef voffset,
-                                      LLVMValueRef soffset,
-                                      unsigned inst_offset,
-                                      unsigned glc,
-                                      unsigned slc)
-{
-       struct gallivm_state *gallivm = &ctx->gallivm;
-       unsigned func = CLAMP(num_channels, 1, 3) - 1;
-
-       if (HAVE_LLVM >= 0x309) {
-               LLVMValueRef args[] = {
-                       LLVMBuildBitCast(gallivm->builder, rsrc, ctx->v4i32, ""),
-                       vindex ? vindex : LLVMConstInt(ctx->i32, 0, 0),
-                       LLVMConstInt(ctx->i32, inst_offset, 0),
-                       LLVMConstInt(ctx->i1, glc, 0),
-                       LLVMConstInt(ctx->i1, slc, 0)
-               };
-
-               LLVMTypeRef types[] = {ctx->f32, LLVMVectorType(ctx->f32, 2),
-                                      ctx->v4f32};
-               const char *type_names[] = {"f32", "v2f32", "v4f32"};
-               char name[256];
-
-               if (voffset) {
-                       args[2] = LLVMBuildAdd(gallivm->builder, args[2], voffset,
-                                              "");
-               }
-
-               if (soffset) {
-                       args[2] = LLVMBuildAdd(gallivm->builder, args[2], soffset,
-                                              "");
-               }
-
-               snprintf(name, sizeof(name), "llvm.amdgcn.buffer.load.%s",
-                        type_names[func]);
-
-               return lp_build_intrinsic(gallivm->builder, name, types[func], args,
-                                         ARRAY_SIZE(args), LP_FUNC_ATTR_READONLY);
-       } else {
-               LLVMValueRef args[] = {
-                       LLVMBuildBitCast(gallivm->builder, rsrc, ctx->v16i8, ""),
-                       voffset ? voffset : vindex,
-                       soffset,
-                       LLVMConstInt(ctx->i32, inst_offset, 0),
-                       LLVMConstInt(ctx->i32, voffset ? 1 : 0, 0), // offen
-                       LLVMConstInt(ctx->i32, vindex ? 1 : 0, 0), //idxen
-                       LLVMConstInt(ctx->i32, glc, 0),
-                       LLVMConstInt(ctx->i32, slc, 0),
-                       LLVMConstInt(ctx->i32, 0, 0), // TFE
-               };
-
-               LLVMTypeRef types[] = {ctx->i32, LLVMVectorType(ctx->i32, 2),
-                                      ctx->v4i32};
-               const char *type_names[] = {"i32", "v2i32", "v4i32"};
-               const char *arg_type = "i32";
-               char name[256];
-
-               if (voffset && vindex) {
-                       LLVMValueRef vaddr[] = {vindex, voffset};
-
-                       arg_type = "v2i32";
-                       args[1] = lp_build_gather_values(gallivm, vaddr, 2);
-               }
-
-               snprintf(name, sizeof(name), "llvm.SI.buffer.load.dword.%s.%s",
-                        type_names[func], arg_type);
-
-               return lp_build_intrinsic(gallivm->builder, name, types[func], args,
-                                         ARRAY_SIZE(args), LP_FUNC_ATTR_READONLY);
-       }
-}
-
 static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
                                 enum tgsi_opcode_type type, unsigned swizzle,
                                 LLVMValueRef buffer, LLVMValueRef offset,
@@ -905,25 +832,25 @@ static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
        LLVMTypeRef vec_type = LLVMVectorType(llvm_type, 4);
 
        if (swizzle == ~0) {
-               value = build_buffer_load(ctx, buffer, 4, NULL, base, offset,
-                                         0, 1, 0);
+               value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
+                                            0, 1, 0);
 
                return LLVMBuildBitCast(gallivm->builder, value, vec_type, "");
        }
 
        if (!tgsi_type_is_64bit(type)) {
-               value = build_buffer_load(ctx, buffer, 4, NULL, base, offset,
-                                         0, 1, 0);
+               value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
+                                            0, 1, 0);
 
                value = LLVMBuildBitCast(gallivm->builder, value, vec_type, "");
                return LLVMBuildExtractElement(gallivm->builder, value,
                                    lp_build_const_int32(gallivm, swizzle), "");
        }
 
-       value = build_buffer_load(ctx, buffer, 1, NULL, base, offset,
+       value = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
                                  swizzle * 4, 1, 0);
 
-       value2 = build_buffer_load(ctx, buffer, 1, NULL, base, offset,
+       value2 = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
                                   swizzle * 4 + 4, 1, 0);
 
        return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
@@ -957,12 +884,12 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
        dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
                            lp_build_const_int32(gallivm, swizzle));
 
-       value = build_indexed_load(ctx, ctx->lds, dw_addr, false);
+       value = ac_build_indexed_load(&ctx->ac, ctx->lds, dw_addr, false);
        if (tgsi_type_is_64bit(type)) {
                LLVMValueRef value2;
                dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
                                       lp_build_const_int32(gallivm, 1));
-               value2 = build_indexed_load(ctx, ctx->lds, dw_addr, false);
+               value2 = ac_build_indexed_load(&ctx->ac, ctx->lds, dw_addr, false);
                return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
        }
 
@@ -988,8 +915,8 @@ static void lds_store(struct lp_build_tgsi_context *bld_base,
                            lp_build_const_int32(gallivm, swizzle));
 
        value = LLVMBuildBitCast(gallivm->builder, value, ctx->i32, "");
-       build_indexed_store(ctx, ctx->lds,
-                           dw_addr, value);
+       ac_build_indexed_store(&ctx->ac, ctx->lds,
+                              dw_addr, value);
 }
 
 static LLVMValueRef fetch_input_tcs(
@@ -1038,7 +965,7 @@ static LLVMValueRef fetch_input_tes(
 
        rw_buffers = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_RW_BUFFERS);
-       buffer = build_indexed_load_const(ctx, rw_buffers,
+       buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
                        lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
 
        base = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
@@ -1080,7 +1007,7 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
 
        rw_buffers = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_RW_BUFFERS);
-       buffer = build_indexed_load_const(ctx, rw_buffers,
+       buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
                        lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
 
        base = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
@@ -1099,17 +1026,17 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
                values[chan_index] = value;
 
                if (inst->Dst[0].Register.WriteMask != 0xF) {
-                       build_tbuffer_store_dwords(ctx, buffer, value, 1,
-                                                  buf_addr, base,
-                                                  4 * chan_index);
+                       ac_build_tbuffer_store_dwords(&ctx->ac, buffer, value, 1,
+                                                     buf_addr, base,
+                                                     4 * chan_index);
                }
        }
 
        if (inst->Dst[0].Register.WriteMask == 0xF) {
                LLVMValueRef value = lp_build_gather_values(bld_base->base.gallivm,
                                                            values, 4);
-               build_tbuffer_store_dwords(ctx, buffer, value, 4, buf_addr,
-                                          base, 0);
+               ac_build_tbuffer_store_dwords(&ctx->ac, buffer, value, 4, buf_addr,
+                                             base, 0);
        }
 }
 
@@ -1122,7 +1049,7 @@ static LLVMValueRef fetch_input_gs(
        struct lp_build_context *base = &bld_base->base;
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct si_shader *shader = ctx->shader;
-       struct lp_build_context *uint = &ctx->soa.bld_base.uint_bld;
+       struct lp_build_context *uint = &ctx->bld_base.uint_bld;
        struct gallivm_state *gallivm = base->gallivm;
        LLVMValueRef vtx_offset;
        LLVMValueRef args[9];
@@ -1221,80 +1148,6 @@ static int lookup_interp_param_index(unsigned interpolate, unsigned location)
        }
 }
 
-static LLVMValueRef build_fs_interp(
-       struct lp_build_tgsi_context *bld_base,
-       LLVMValueRef llvm_chan,
-       LLVMValueRef attr_number,
-       LLVMValueRef params,
-       LLVMValueRef i,
-       LLVMValueRef j) {
-
-       struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef args[5];
-       LLVMValueRef p1;
-       if (HAVE_LLVM < 0x0400) {
-               LLVMValueRef ij[2];
-               ij[0] = LLVMBuildBitCast(gallivm->builder, i, ctx->i32, "");
-               ij[1] = LLVMBuildBitCast(gallivm->builder, j, ctx->i32, "");
-
-               args[0] = llvm_chan;
-               args[1] = attr_number;
-               args[2] = params;
-               args[3] = lp_build_gather_values(gallivm, ij, 2);
-               return lp_build_intrinsic(gallivm->builder, "llvm.SI.fs.interp",
-                                         ctx->f32, args, 4,
-                                         LP_FUNC_ATTR_READNONE);
-       }
-
-       args[0] = i;
-       args[1] = llvm_chan;
-       args[2] = attr_number;
-       args[3] = params;
-
-       p1 = lp_build_intrinsic(gallivm->builder, "llvm.amdgcn.interp.p1",
-                               ctx->f32, args, 4, LP_FUNC_ATTR_READNONE);
-
-       args[0] = p1;
-       args[1] = j;
-       args[2] = llvm_chan;
-       args[3] = attr_number;
-       args[4] = params;
-
-       return lp_build_intrinsic(gallivm->builder, "llvm.amdgcn.interp.p2",
-                                 ctx->f32, args, 5, LP_FUNC_ATTR_READNONE);
-}
-
-static LLVMValueRef build_fs_interp_mov(
-       struct lp_build_tgsi_context *bld_base,
-       LLVMValueRef parameter,
-       LLVMValueRef llvm_chan,
-       LLVMValueRef attr_number,
-       LLVMValueRef params) {
-
-       struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef args[4];
-       if (HAVE_LLVM < 0x0400) {
-               args[0] = llvm_chan;
-               args[1] = attr_number;
-               args[2] = params;
-
-               return lp_build_intrinsic(gallivm->builder,
-                                         "llvm.SI.fs.constant",
-                                         ctx->f32, args, 3,
-                                         LP_FUNC_ATTR_READNONE);
-       }
-
-       args[0] = parameter;
-       args[1] = llvm_chan;
-       args[2] = attr_number;
-       args[3] = params;
-
-       return lp_build_intrinsic(gallivm->builder, "llvm.amdgcn.interp.mov",
-                                 ctx->f32, args, 4, LP_FUNC_ATTR_READNONE);
-}
-
 /**
  * Interpolate a fragment shader input.
  *
@@ -1320,7 +1173,7 @@ static void interp_fs_input(struct si_shader_context *ctx,
                            LLVMValueRef face,
                            LLVMValueRef result[4])
 {
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
        struct lp_build_context *base = &bld_base->base;
        struct lp_build_context *uint = &bld_base->uint_bld;
        struct gallivm_state *gallivm = base->gallivm;
@@ -1378,17 +1231,17 @@ static void interp_fs_input(struct si_shader_context *ctx,
                        LLVMValueRef front, back;
 
                        if (interp) {
-                               front = build_fs_interp(bld_base, llvm_chan,
+                               front = ac_build_fs_interp(&ctx->ac, llvm_chan,
                                                        attr_number, prim_mask,
                                                        i, j);
-                               back = build_fs_interp(bld_base, llvm_chan,
+                               back = ac_build_fs_interp(&ctx->ac, llvm_chan,
                                                        back_attr_number, prim_mask,
                                                        i, j);
                        } else {
-                               front = build_fs_interp_mov(bld_base,
+                               front = ac_build_fs_interp_mov(&ctx->ac,
                                        lp_build_const_int32(gallivm, 2), /* P0 */
                                        llvm_chan, attr_number, prim_mask);
-                               back = build_fs_interp_mov(bld_base,
+                               back = ac_build_fs_interp_mov(&ctx->ac,
                                        lp_build_const_int32(gallivm, 2), /* P0 */
                                        llvm_chan, back_attr_number, prim_mask);
                        }
@@ -1401,12 +1254,12 @@ static void interp_fs_input(struct si_shader_context *ctx,
                }
        } else if (semantic_name == TGSI_SEMANTIC_FOG) {
                if (interp) {
-                       result[0] = build_fs_interp(bld_base, uint->zero,
-                                               attr_number, prim_mask, i, j);
+                       result[0] = ac_build_fs_interp(&ctx->ac, uint->zero,
+                                                      attr_number, prim_mask, i, j);
                } else {
-                       result[0] = build_fs_interp_mov(bld_base, uint->zero,
-                               lp_build_const_int32(gallivm, 2), /* P0 */
-                               attr_number, prim_mask);
+                       result[0] = ac_build_fs_interp_mov(&ctx->ac, uint->zero,
+                                                          lp_build_const_int32(gallivm, 2), /* P0 */
+                                                          attr_number, prim_mask);
                }
                result[1] =
                result[2] = lp_build_const_float(gallivm, 0.0f);
@@ -1416,10 +1269,10 @@ static void interp_fs_input(struct si_shader_context *ctx,
                        LLVMValueRef llvm_chan = lp_build_const_int32(gallivm, chan);
 
                        if (interp) {
-                               result[chan] = build_fs_interp(bld_base,
+                               result[chan] = ac_build_fs_interp(&ctx->ac,
                                        llvm_chan, attr_number, prim_mask, i, j);
                        } else {
-                               result[chan] = build_fs_interp_mov(bld_base,
+                               result[chan] = ac_build_fs_interp_mov(&ctx->ac,
                                        lp_build_const_int32(gallivm, 2), /* P0 */
                                        llvm_chan, attr_number, prim_mask);
                        }
@@ -1433,9 +1286,9 @@ static void declare_input_fs(
        const struct tgsi_full_declaration *decl,
        LLVMValueRef out[4])
 {
-       struct lp_build_context *base = &radeon_bld->soa.bld_base.base;
+       struct lp_build_context *base = &radeon_bld->bld_base.base;
        struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->soa.bld_base);
+               si_shader_context(&radeon_bld->bld_base);
        struct si_shader *shader = ctx->shader;
        LLVMValueRef main_fn = radeon_bld->main_fn;
        LLVMValueRef interp_param = NULL;
@@ -1479,53 +1332,10 @@ static void declare_input_fs(
 
 static LLVMValueRef get_sample_id(struct si_shader_context *radeon_bld)
 {
-       return unpack_param(si_shader_context(&radeon_bld->soa.bld_base),
+       return unpack_param(si_shader_context(&radeon_bld->bld_base),
                            SI_PARAM_ANCILLARY, 8, 4);
 }
 
-/**
- * Set range metadata on an instruction.  This can only be used on load and
- * call instructions.  If you know an instruction can only produce the values
- * 0, 1, 2, you would do set_range_metadata(value, 0, 3);
- * \p lo is the minimum value inclusive.
- * \p hi is the maximum value exclusive.
- */
-static void set_range_metadata(struct si_shader_context *ctx,
-                              LLVMValueRef value, unsigned lo, unsigned hi)
-{
-       LLVMValueRef range_md, md_args[2];
-       LLVMTypeRef type = LLVMTypeOf(value);
-       LLVMContextRef context = LLVMGetTypeContext(type);
-
-       md_args[0] = LLVMConstInt(type, lo, false);
-       md_args[1] = LLVMConstInt(type, hi, false);
-       range_md = LLVMMDNodeInContext(context, md_args, 2);
-       LLVMSetMetadata(value, ctx->range_md_kind, range_md);
-}
-
-static LLVMValueRef get_thread_id(struct si_shader_context *ctx)
-{
-       struct gallivm_state *gallivm = &ctx->gallivm;
-       LLVMValueRef tid;
-
-       if (HAVE_LLVM < 0x0308) {
-               tid = lp_build_intrinsic(gallivm->builder, "llvm.SI.tid",
-                               ctx->i32,   NULL, 0, LP_FUNC_ATTR_READNONE);
-       } else {
-               LLVMValueRef tid_args[2];
-               tid_args[0] = lp_build_const_int32(gallivm, 0xffffffff);
-               tid_args[1] = lp_build_const_int32(gallivm, 0);
-               tid_args[1] = lp_build_intrinsic(gallivm->builder,
-                                       "llvm.amdgcn.mbcnt.lo", ctx->i32,
-                                       tid_args, 2, LP_FUNC_ATTR_READNONE);
-
-               tid = lp_build_intrinsic(gallivm->builder,
-                                       "llvm.amdgcn.mbcnt.hi", ctx->i32,
-                                       tid_args, 2, LP_FUNC_ATTR_READNONE);
-       }
-       set_range_metadata(ctx, tid, 0, 64);
-       return tid;
-}
 
 /**
  * Load a dword from a constant buffer.
@@ -1544,13 +1354,13 @@ static LLVMValueRef buffer_load_const(struct si_shader_context *ctx,
 static LLVMValueRef load_sample_position(struct si_shader_context *radeon_bld, LLVMValueRef sample_id)
 {
        struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->soa.bld_base);
-       struct lp_build_context *uint_bld = &radeon_bld->soa.bld_base.uint_bld;
+               si_shader_context(&radeon_bld->bld_base);
+       struct lp_build_context *uint_bld = &radeon_bld->bld_base.uint_bld;
        struct gallivm_state *gallivm = &radeon_bld->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        LLVMValueRef desc = LLVMGetParam(ctx->main_fn, SI_PARAM_RW_BUFFERS);
        LLVMValueRef buf_index = lp_build_const_int32(gallivm, SI_PS_CONST_SAMPLE_POSITIONS);
-       LLVMValueRef resource = build_indexed_load_const(ctx, desc, buf_index);
+       LLVMValueRef resource = ac_build_indexed_load_const(&ctx->ac, desc, buf_index);
 
        /* offset = sample_id * 8  (8 = 2 floats containing samplepos.xy) */
        LLVMValueRef offset0 = lp_build_mul_imm(uint_bld, sample_id, 8);
@@ -1572,8 +1382,8 @@ static void declare_system_value(
        const struct tgsi_full_declaration *decl)
 {
        struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->soa.bld_base);
-       struct lp_build_context *bld = &radeon_bld->soa.bld_base.base;
+               si_shader_context(&radeon_bld->bld_base);
+       struct lp_build_context *bld = &radeon_bld->bld_base.base;
        struct gallivm_state *gallivm = &radeon_bld->gallivm;
        LLVMValueRef value = 0;
 
@@ -1627,7 +1437,7 @@ static void declare_system_value(
                        LLVMGetParam(radeon_bld->main_fn, SI_PARAM_POS_X_FLOAT),
                        LLVMGetParam(radeon_bld->main_fn, SI_PARAM_POS_Y_FLOAT),
                        LLVMGetParam(radeon_bld->main_fn, SI_PARAM_POS_Z_FLOAT),
-                       lp_build_emit_llvm_unary(&radeon_bld->soa.bld_base, TGSI_OPCODE_RCP,
+                       lp_build_emit_llvm_unary(&radeon_bld->bld_base, TGSI_OPCODE_RCP,
                                                 LLVMGetParam(radeon_bld->main_fn,
                                                              SI_PARAM_POS_W_FLOAT)),
                };
@@ -1650,9 +1460,9 @@ static void declare_system_value(
                        lp_build_const_float(gallivm, 0),
                        lp_build_const_float(gallivm, 0)
                };
-               pos[0] = lp_build_emit_llvm_unary(&radeon_bld->soa.bld_base,
+               pos[0] = lp_build_emit_llvm_unary(&radeon_bld->bld_base,
                                                  TGSI_OPCODE_FRC, pos[0]);
-               pos[1] = lp_build_emit_llvm_unary(&radeon_bld->soa.bld_base,
+               pos[1] = lp_build_emit_llvm_unary(&radeon_bld->bld_base,
                                                  TGSI_OPCODE_FRC, pos[1]);
                value = lp_build_gather_values(gallivm, pos, 4);
                break;
@@ -1701,14 +1511,14 @@ static void declare_system_value(
 
                rw_buffers = LLVMGetParam(ctx->main_fn,
                                        SI_PARAM_RW_BUFFERS);
-               buffer = build_indexed_load_const(ctx, rw_buffers,
+               buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
                        lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
 
                base = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
                addr = get_tcs_tes_buffer_address(ctx, NULL,
                                          lp_build_const_int32(gallivm, param));
 
-               value = buffer_load(&radeon_bld->soa.bld_base, TGSI_TYPE_FLOAT,
+               value = buffer_load(&radeon_bld->bld_base, TGSI_TYPE_FLOAT,
                                    ~0, buffer, base, addr);
 
                break;
@@ -1722,7 +1532,7 @@ static void declare_system_value(
 
                slot = lp_build_const_int32(gallivm, SI_HS_CONST_DEFAULT_TESS_LEVELS);
                buf = LLVMGetParam(ctx->main_fn, SI_PARAM_RW_BUFFERS);
-               buf = build_indexed_load_const(ctx, buf, slot);
+               buf = ac_build_indexed_load_const(&ctx->ac, buf, slot);
                offset = decl->Semantic.Name == TGSI_SEMANTIC_DEFAULT_TESSINNER_SI ? 4 : 0;
 
                for (i = 0; i < 4; i++)
@@ -1733,7 +1543,7 @@ static void declare_system_value(
        }
 
        case TGSI_SEMANTIC_PRIMID:
-               value = get_primitive_id(&radeon_bld->soa.bld_base, 0);
+               value = get_primitive_id(&radeon_bld->bld_base, 0);
                break;
 
        case TGSI_SEMANTIC_GRID_SIZE:
@@ -1771,16 +1581,19 @@ static void declare_system_value(
                value = LLVMGetParam(radeon_bld->main_fn, SI_PARAM_THREAD_ID);
                break;
 
-#if HAVE_LLVM >= 0x0309
        case TGSI_SEMANTIC_HELPER_INVOCATION:
-               value = lp_build_intrinsic(gallivm->builder,
-                                          "llvm.amdgcn.ps.live",
-                                          ctx->i1, NULL, 0,
-                                          LP_FUNC_ATTR_READNONE);
-               value = LLVMBuildNot(gallivm->builder, value, "");
-               value = LLVMBuildSExt(gallivm->builder, value, ctx->i32, "");
+               if (HAVE_LLVM >= 0x0309) {
+                       value = lp_build_intrinsic(gallivm->builder,
+                                                  "llvm.amdgcn.ps.live",
+                                                  ctx->i1, NULL, 0,
+                                                  LP_FUNC_ATTR_READNONE);
+                       value = LLVMBuildNot(gallivm->builder, value, "");
+                       value = LLVMBuildSExt(gallivm->builder, value, ctx->i32, "");
+               } else {
+                       assert(!"TGSI_SEMANTIC_HELPER_INVOCATION unsupported");
+                       return;
+               }
                break;
-#endif
 
        default:
                assert(!"unknown system value");
@@ -1794,7 +1607,7 @@ static void declare_compute_memory(struct si_shader_context *radeon_bld,
                                    const struct tgsi_full_declaration *decl)
 {
        struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->soa.bld_base);
+               si_shader_context(&radeon_bld->bld_base);
        struct si_shader_selector *sel = ctx->shader->selector;
        struct gallivm_state *gallivm = &radeon_bld->gallivm;
 
@@ -1819,7 +1632,7 @@ static LLVMValueRef load_const_buffer_desc(struct si_shader_context *ctx, int i)
        LLVMValueRef list_ptr = LLVMGetParam(ctx->main_fn,
                                             SI_PARAM_CONST_BUFFERS);
 
-       return build_indexed_load_const(ctx, list_ptr,
+       return ac_build_indexed_load_const(&ctx->ac, list_ptr,
                                        LLVMConstInt(ctx->i32, i, 0));
 }
 
@@ -1855,12 +1668,12 @@ static LLVMValueRef fetch_constant(
                index = get_bounded_indirect_index(ctx, &reg->DimIndirect,
                                                   reg->Dimension.Index,
                                                   SI_NUM_CONST_BUFFERS);
-               bufp = build_indexed_load_const(ctx, ptr, index);
+               bufp = ac_build_indexed_load_const(&ctx->ac, ptr, index);
        } else
                bufp = load_const_buffer_desc(ctx, buf);
 
        if (reg->Register.Indirect) {
-               addr = ctx->soa.addr[ireg->Index][ireg->Swizzle];
+               addr = ctx->addrs[ireg->Index][ireg->Swizzle];
                addr = LLVMBuildLoad(base->gallivm->builder, addr, "load addr reg");
                addr = lp_build_mul_imm(&bld_base->uint_bld, addr, 16);
                addr = lp_build_add(&bld_base->uint_bld, addr,
@@ -1915,8 +1728,7 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
                                     LLVMValueRef *args)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct lp_build_context *uint =
-                               &ctx->soa.bld_base.uint_bld;
+       struct lp_build_context *uint = &ctx->bld_base.uint_bld;
        struct lp_build_context *base = &bld_base->base;
        struct gallivm_state *gallivm = base->gallivm;
        LLVMBuilderRef builder = base->gallivm->builder;
@@ -2151,7 +1963,7 @@ static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base,
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct lp_build_context *base = &bld_base->base;
-       struct lp_build_context *uint = &ctx->soa.bld_base.uint_bld;
+       struct lp_build_context *uint = &ctx->bld_base.uint_bld;
        unsigned reg_index;
        unsigned chan;
        unsigned const_chan;
@@ -2159,7 +1971,7 @@ static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base,
        LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, SI_PARAM_RW_BUFFERS);
        LLVMValueRef constbuf_index = lp_build_const_int32(base->gallivm,
                                                           SI_VS_CONST_CLIP_PLANES);
-       LLVMValueRef const_resource = build_indexed_load_const(ctx, ptr, constbuf_index);
+       LLVMValueRef const_resource = ac_build_indexed_load_const(&ctx->ac, ptr, constbuf_index);
 
        for (reg_index = 0; reg_index < 2; reg_index ++) {
                LLVMValueRef *args = pos[2 + reg_index];
@@ -2258,11 +2070,11 @@ static void emit_streamout_output(struct si_shader_context *ctx,
                break;
        }
 
-       build_tbuffer_store_dwords(ctx, so_buffers[buf_idx],
-                                  vdata, num_comps,
-                                  so_write_offsets[buf_idx],
-                                  LLVMConstInt(ctx->i32, 0, 0),
-                                  stream_out->dst_offset * 4);
+       ac_build_tbuffer_store_dwords(&ctx->ac, so_buffers[buf_idx],
+                                     vdata, num_comps,
+                                     so_write_offsets[buf_idx],
+                                     LLVMConstInt(ctx->i32, 0, 0),
+                                     stream_out->dst_offset * 4);
 }
 
 /**
@@ -2284,7 +2096,7 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
        LLVMValueRef so_vtx_count =
                unpack_param(ctx, ctx->param_streamout_config, 16, 7);
 
-       LLVMValueRef tid = get_thread_id(ctx);
+       LLVMValueRef tid = ac_get_thread_id(&ctx->ac);
 
        /* can_emit = tid < so_vtx_count; */
        LLVMValueRef can_emit =
@@ -2322,7 +2134,7 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
                        LLVMValueRef offset = lp_build_const_int32(gallivm,
                                                                   SI_VS_STREAMOUT_BUF0 + i);
 
-                       so_buffers[i] = build_indexed_load_const(ctx, buf_ptr, offset);
+                       so_buffers[i] = ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
 
                        LLVMValueRef so_offset = LLVMGetParam(ctx->main_fn,
                                                              ctx->param_streamout_offset[i]);
@@ -2359,8 +2171,7 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base,
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct si_shader *shader = ctx->shader;
        struct lp_build_context *base = &bld_base->base;
-       struct lp_build_context *uint =
-                               &ctx->soa.bld_base.uint_bld;
+       struct lp_build_context *uint = &ctx->bld_base.uint_bld;
        LLVMValueRef args[9];
        LLVMValueRef pos_args[4][9] = { { 0 } };
        LLVMValueRef psize_value = NULL, edgeflag_value = NULL, layer_value = NULL, viewport_index_value = NULL;
@@ -2420,15 +2231,6 @@ handle_semantic:
                case TGSI_SEMANTIC_POSITION:
                        target = V_008DFC_SQ_EXP_POS;
                        break;
-               case TGSI_SEMANTIC_COLOR:
-               case TGSI_SEMANTIC_BCOLOR:
-                       if (!export_param)
-                               continue;
-                       target = V_008DFC_SQ_EXP_PARAM + param_count;
-                       assert(i < ARRAY_SIZE(shader->info.vs_output_param_offset));
-                       shader->info.vs_output_param_offset[i] = param_count;
-                       param_count++;
-                       break;
                case TGSI_SEMANTIC_CLIPDIST:
                        if (shader->key.opt.hw_vs.clip_disable) {
                                semantic_name = TGSI_SEMANTIC_GENERIC;
@@ -2441,6 +2243,8 @@ handle_semantic:
                                continue;
                        si_llvm_emit_clipvertex(bld_base, pos_args, outputs[i].values);
                        continue;
+               case TGSI_SEMANTIC_COLOR:
+               case TGSI_SEMANTIC_BCOLOR:
                case TGSI_SEMANTIC_PRIMID:
                case TGSI_SEMANTIC_FOG:
                case TGSI_SEMANTIC_TEXCOORD:
@@ -2573,7 +2377,7 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
        invocation_id = unpack_param(ctx, SI_PARAM_REL_IDS, 8, 5);
 
        rw_buffers = LLVMGetParam(ctx->main_fn, SI_PARAM_RW_BUFFERS);
-       buffer = build_indexed_load_const(ctx, rw_buffers,
+       buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
                        lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
 
        buffer_offset = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
@@ -2599,8 +2403,8 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
                LLVMValueRef value = lds_load(bld_base, TGSI_TYPE_SIGNED, ~0,
                                              lds_ptr);
 
-               build_tbuffer_store_dwords(ctx, buffer, value, 4, buffer_addr,
-                                          buffer_offset, 0);
+               ac_build_tbuffer_store_dwords(&ctx->ac, buffer, value, 4, buffer_addr,
+                                             buffer_offset, 0);
        }
 }
 
@@ -2689,7 +2493,7 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
        /* Get the buffer. */
        rw_buffers = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_RW_BUFFERS);
-       buffer = build_indexed_load_const(ctx, rw_buffers,
+       buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
                        lp_build_const_int32(gallivm, SI_HS_RING_TESS_FACTOR));
 
        /* Get the offset. */
@@ -2703,18 +2507,18 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
                                  rel_patch_id, bld_base->uint_bld.zero, ""));
 
        /* Store the dynamic HS control word. */
-       build_tbuffer_store_dwords(ctx, buffer,
-                                  lp_build_const_int32(gallivm, 0x80000000),
-                                  1, lp_build_const_int32(gallivm, 0), tf_base, 0);
+       ac_build_tbuffer_store_dwords(&ctx->ac, buffer,
+                                     lp_build_const_int32(gallivm, 0x80000000),
+                                     1, lp_build_const_int32(gallivm, 0), tf_base, 0);
 
        lp_build_endif(&inner_if_ctx);
 
        /* Store the tessellation factors. */
-       build_tbuffer_store_dwords(ctx, buffer, vec0,
-                                  MIN2(stride, 4), byteoffset, tf_base, 4);
+       ac_build_tbuffer_store_dwords(&ctx->ac, buffer, vec0,
+                                     MIN2(stride, 4), byteoffset, tf_base, 4);
        if (vec1)
-               build_tbuffer_store_dwords(ctx, buffer, vec1,
-                                          stride - 4, byteoffset, tf_base, 20);
+               ac_build_tbuffer_store_dwords(&ctx->ac, buffer, vec1,
+                                             stride - 4, byteoffset, tf_base, 20);
        lp_build_endif(&if_ctx);
 }
 
@@ -2783,7 +2587,7 @@ static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base)
        /* Write outputs to LDS. The next shader (TCS aka HS) will read
         * its inputs from it. */
        for (i = 0; i < info->num_outputs; i++) {
-               LLVMValueRef *out_ptr = ctx->soa.outputs[i];
+               LLVMValueRef *out_ptr = ctx->outputs[i];
                unsigned name = info->output_semantic_name[i];
                unsigned index = info->output_semantic_index[i];
                int param = si_shader_io_get_unique_index(name, index);
@@ -2809,8 +2613,7 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base)
        int i;
 
        for (i = 0; i < info->num_outputs; i++) {
-               LLVMValueRef *out_ptr =
-                       ctx->soa.outputs[i];
+               LLVMValueRef *out_ptr = ctx->outputs[i];
                int param_index;
 
                if (info->output_semantic_name[i] == TGSI_SEMANTIC_VIEWPORT_INDEX ||
@@ -2824,14 +2627,14 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base)
                        LLVMValueRef out_val = LLVMBuildLoad(gallivm->builder, out_ptr[chan], "");
                        out_val = LLVMBuildBitCast(gallivm->builder, out_val, ctx->i32, "");
 
-                       build_tbuffer_store(ctx,
-                                           ctx->esgs_ring,
-                                           out_val, 1,
-                                           LLVMGetUndef(ctx->i32), soffset,
-                                           (4 * param_index + chan) * 4,
-                                           V_008F0C_BUF_DATA_FORMAT_32,
-                                           V_008F0C_BUF_NUM_FORMAT_UINT,
-                                           0, 0, 1, 1, 0);
+                       ac_build_tbuffer_store(&ctx->ac,
+                                              ctx->esgs_ring,
+                                              out_val, 1,
+                                              LLVMGetUndef(ctx->i32), soffset,
+                                              (4 * param_index + chan) * 4,
+                                              V_008F0C_BUF_DATA_FORMAT_32,
+                                              V_008F0C_BUF_NUM_FORMAT_UINT,
+                                              0, 0, 1, 1, 0);
                }
        }
 }
@@ -2839,13 +2642,9 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base)
 static void si_llvm_emit_gs_epilogue(struct lp_build_tgsi_context *bld_base)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef args[2];
 
-       args[0] = lp_build_const_int32(gallivm, SENDMSG_GS_OP_NOP | SENDMSG_GS_DONE);
-       args[1] = LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID);
-       lp_build_intrinsic(gallivm->builder, "llvm.SI.sendmsg",
-                          ctx->voidt, args, 2, 0);
+       ac_emit_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_NOP | AC_SENDMSG_GS_DONE,
+                       LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
 }
 
 static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
@@ -2887,7 +2686,7 @@ static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
                        }
 
                        for (j = 0; j < 4; j++) {
-                               addr = ctx->soa.outputs[i][j];
+                               addr = ctx->outputs[i][j];
                                val = LLVMBuildLoad(gallivm->builder, addr, "");
                                val = si_llvm_saturate(bld_base, val);
                                LLVMBuildStore(gallivm->builder, val, addr);
@@ -2905,7 +2704,7 @@ static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
                for (j = 0; j < 4; j++) {
                        outputs[i].values[j] =
                                LLVMBuildLoad(gallivm->builder,
-                                             ctx->soa.outputs[i][j],
+                                             ctx->outputs[i][j],
                                              "");
                        outputs[i].vertex_stream[j] =
                                (info->output_streams[i] >> (2 * j)) & 3;
@@ -3155,22 +2954,22 @@ static void si_llvm_return_fs_outputs(struct lp_build_tgsi_context *bld_base)
                case TGSI_SEMANTIC_COLOR:
                        assert(semantic_index < 8);
                        for (j = 0; j < 4; j++) {
-                               LLVMValueRef ptr = ctx->soa.outputs[i][j];
+                               LLVMValueRef ptr = ctx->outputs[i][j];
                                LLVMValueRef result = LLVMBuildLoad(builder, ptr, "");
                                color[semantic_index][j] = result;
                        }
                        break;
                case TGSI_SEMANTIC_POSITION:
                        depth = LLVMBuildLoad(builder,
-                                             ctx->soa.outputs[i][2], "");
+                                             ctx->outputs[i][2], "");
                        break;
                case TGSI_SEMANTIC_STENCIL:
                        stencil = LLVMBuildLoad(builder,
-                                               ctx->soa.outputs[i][1], "");
+                                               ctx->outputs[i][1], "");
                        break;
                case TGSI_SEMANTIC_SAMPLEMASK:
                        samplemask = LLVMBuildLoad(builder,
-                                                  ctx->soa.outputs[i][0], "");
+                                                  ctx->outputs[i][0], "");
                        break;
                default:
                        fprintf(stderr, "Warning: SI unhandled fs output type:%d\n",
@@ -3361,7 +3160,7 @@ shader_buffer_fetch_rsrc(struct si_shader_context *ctx,
                                                   reg->Register.Index,
                                                   SI_NUM_SHADER_BUFFERS);
 
-       return build_indexed_load_const(ctx, rsrc_ptr, index);
+       return ac_build_indexed_load_const(&ctx->ac, rsrc_ptr, index);
 }
 
 static bool tgsi_is_array_sampler(unsigned target)
@@ -3469,11 +3268,11 @@ image_fetch_rsrc(
                                     LLVMConstInt(ctx->i32, 2, 0), "");
                index = LLVMBuildAdd(builder, index,
                                     LLVMConstInt(ctx->i32, 1, 0), "");
-               *rsrc = build_indexed_load_const(ctx, rsrc_ptr, index);
+               *rsrc = ac_build_indexed_load_const(&ctx->ac, rsrc_ptr, index);
                return;
        }
 
-       tmp = build_indexed_load_const(ctx, rsrc_ptr, index);
+       tmp = ac_build_indexed_load_const(&ctx->ac, rsrc_ptr, index);
        if (dcc_off)
                tmp = force_dcc_off(ctx, tmp);
        *rsrc = tmp;
@@ -3666,7 +3465,7 @@ static LLVMValueRef get_memory_ptr(struct si_shader_context *ctx,
        LLVMValueRef offset, ptr;
        int addr_space;
 
-       offset = lp_build_emit_fetch(&ctx->soa.bld_base, inst, arg, 0);
+       offset = lp_build_emit_fetch(&ctx->bld_base, inst, arg, 0);
        offset = LLVMBuildBitCast(builder, offset, ctx->i32, "");
 
        ptr = ctx->shared_memory;
@@ -3682,7 +3481,7 @@ static void load_emit_memory(
                struct lp_build_emit_data *emit_data)
 {
        const struct tgsi_full_instruction *inst = emit_data->inst;
-       struct lp_build_context *base = &ctx->soa.bld_base.base;
+       struct lp_build_context *base = &ctx->bld_base.base;
        struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        unsigned writemask = inst->Dst[0].Register.WriteMask;
@@ -3848,7 +3647,7 @@ static void store_emit_buffer(
        const struct tgsi_full_instruction *inst = emit_data->inst;
        struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
-       struct lp_build_context *uint_bld = &ctx->soa.bld_base.uint_bld;
+       struct lp_build_context *uint_bld = &ctx->bld_base.uint_bld;
        LLVMValueRef base_data = emit_data->args[0];
        LLVMValueRef base_offset = emit_data->args[3];
        unsigned writemask = inst->Dst[0].Register.WriteMask;
@@ -3919,7 +3718,7 @@ static void store_emit_memory(
 {
        const struct tgsi_full_instruction *inst = emit_data->inst;
        struct gallivm_state *gallivm = &ctx->gallivm;
-       struct lp_build_context *base = &ctx->soa.bld_base.base;
+       struct lp_build_context *base = &ctx->bld_base.base;
        LLVMBuilderRef builder = gallivm->builder;
        unsigned writemask = inst->Dst[0].Register.WriteMask;
        LLVMValueRef ptr, derived_ptr, data, index;
@@ -3931,7 +3730,7 @@ static void store_emit_memory(
                if (!(writemask & (1 << chan))) {
                        continue;
                }
-               data = lp_build_emit_fetch(&ctx->soa.bld_base, inst, 1, chan);
+               data = lp_build_emit_fetch(&ctx->bld_base, inst, 1, chan);
                index = lp_build_const_int32(gallivm, chan);
                derived_ptr = LLVMBuildGEP(builder, ptr, &index, 1, "");
                LLVMBuildStore(builder, data, derived_ptr);
@@ -4049,12 +3848,12 @@ static void atomic_emit_memory(struct si_shader_context *ctx,
 
        ptr = get_memory_ptr(ctx, inst, ctx->i32, 1);
 
-       arg = lp_build_emit_fetch(&ctx->soa.bld_base, inst, 2, 0);
+       arg = lp_build_emit_fetch(&ctx->bld_base, inst, 2, 0);
        arg = LLVMBuildBitCast(builder, arg, ctx->i32, "");
 
        if (inst->Instruction.Opcode == TGSI_OPCODE_ATOMCAS) {
                LLVMValueRef new_data;
-               new_data = lp_build_emit_fetch(&ctx->soa.bld_base,
+               new_data = lp_build_emit_fetch(&ctx->bld_base,
                                               inst, 3, 0);
 
                new_data = LLVMBuildBitCast(builder, new_data, ctx->i32, "");
@@ -4313,7 +4112,7 @@ static LLVMValueRef load_sampler_desc_custom(struct si_shader_context *ctx,
                break;
        }
 
-       return build_indexed_load_const(ctx, list, index);
+       return ac_build_indexed_load_const(&ctx->ac, list, index);
 }
 
 static LLVMValueRef load_sampler_desc(struct si_shader_context *ctx,
@@ -4610,7 +4409,11 @@ static void tex_fetch_args(
            target == TGSI_TEXTURE_CUBE_ARRAY ||
            target == TGSI_TEXTURE_SHADOWCUBE ||
            target == TGSI_TEXTURE_SHADOWCUBE_ARRAY)
-               si_prepare_cube_coords(bld_base, emit_data, coords, derivs);
+               ac_prepare_cube_coords(&ctx->ac,
+                                      opcode == TGSI_OPCODE_TXD,
+                                      target == TGSI_TEXTURE_CUBE_ARRAY ||
+                                      target == TGSI_TEXTURE_SHADOWCUBE_ARRAY,
+                                      coords, derivs);
 
        if (opcode == TGSI_OPCODE_TXD)
                for (int i = 0; i < num_deriv_channels * 2; i++)
@@ -4658,16 +4461,12 @@ static void tex_fetch_args(
                struct lp_build_context *uint_bld = &bld_base->uint_bld;
                struct lp_build_emit_data txf_emit_data = *emit_data;
                LLVMValueRef txf_address[4];
-               unsigned txf_count = count;
+               /* We only need .xy for non-arrays, and .xyz for arrays. */
+               unsigned txf_count = target == TGSI_TEXTURE_2D_MSAA ? 2 : 3;
                struct tgsi_full_instruction inst = {};
 
                memcpy(txf_address, address, sizeof(txf_address));
 
-               if (target == TGSI_TEXTURE_2D_MSAA) {
-                       txf_address[2] = bld_base->uint_bld.zero;
-               }
-               txf_address[3] = bld_base->uint_bld.zero;
-
                /* Read FMASK using TXF. */
                inst.Instruction.Opcode = TGSI_OPCODE_TXF;
                inst.Texture.Texture = target;
@@ -4688,7 +4487,7 @@ static void tex_fetch_args(
                                                txf_emit_data.output[0],
                                                uint_bld->zero, "");
 
-               unsigned sample_chan = target == TGSI_TEXTURE_2D_MSAA ? 2 : 3;
+               unsigned sample_chan = txf_count; /* the sample index is last */
 
                LLVMValueRef sample_index4 =
                        LLVMBuildMul(gallivm->builder, address[sample_chan], four, "");
@@ -4724,7 +4523,6 @@ static void tex_fetch_args(
                /* add tex offsets */
                if (inst->Texture.NumOffsets) {
                        struct lp_build_context *uint_bld = &bld_base->uint_bld;
-                       struct lp_build_tgsi_soa_context *bld = lp_soa_context(bld_base);
                        const struct tgsi_texture_offset *off = inst->TexOffsets;
 
                        assert(inst->Texture.NumOffsets == 1);
@@ -4732,7 +4530,7 @@ static void tex_fetch_args(
                        switch (target) {
                        case TGSI_TEXTURE_3D:
                                address[2] = lp_build_add(uint_bld, address[2],
-                                               bld->immediates[off->Index][off->SwizzleZ]);
+                                               ctx->imms[off->Index * TGSI_NUM_CHANNELS + off->SwizzleZ]);
                                /* fall through */
                        case TGSI_TEXTURE_2D:
                        case TGSI_TEXTURE_SHADOW2D:
@@ -4742,7 +4540,7 @@ static void tex_fetch_args(
                        case TGSI_TEXTURE_SHADOW2D_ARRAY:
                                address[1] =
                                        lp_build_add(uint_bld, address[1],
-                                               bld->immediates[off->Index][off->SwizzleY]);
+                                               ctx->imms[off->Index * TGSI_NUM_CHANNELS + off->SwizzleY]);
                                /* fall through */
                        case TGSI_TEXTURE_1D:
                        case TGSI_TEXTURE_SHADOW1D:
@@ -4750,7 +4548,7 @@ static void tex_fetch_args(
                        case TGSI_TEXTURE_SHADOW1D_ARRAY:
                                address[0] =
                                        lp_build_add(uint_bld, address[0],
-                                               bld->immediates[off->Index][off->SwizzleX]);
+                                               ctx->imms[off->Index * TGSI_NUM_CHANNELS + off->SwizzleX]);
                                break;
                                /* texture offsets do not apply to other texture targets */
                        }
@@ -4770,13 +4568,12 @@ static void tex_fetch_args(
 
                /* Get the component index from src1.x for Gather4. */
                if (!tgsi_is_shadow_target(target)) {
-                       LLVMValueRef (*imms)[4] = lp_soa_context(bld_base)->immediates;
                        LLVMValueRef comp_imm;
                        struct tgsi_src_register src1 = inst->Src[1].Register;
 
                        assert(src1.File == TGSI_FILE_IMMEDIATE);
 
-                       comp_imm = imms[src1.Index][src1.SwizzleX];
+                       comp_imm = ctx->imms[src1.Index * TGSI_NUM_CHANNELS + src1.SwizzleX];
                        gather_comp = LLVMConstIntGetZExtValue(comp_imm);
                        gather_comp = CLAMP(gather_comp, 0, 3);
                }
@@ -4821,9 +4618,9 @@ static void si_lower_gather4_integer(struct si_shader_context *ctx,
                set_tex_fetch_args(ctx, &txq_emit_data, TGSI_OPCODE_TXQ,
                                   txq_inst.Texture.Texture,
                                   emit_data->args[1], NULL,
-                                  &ctx->soa.bld_base.uint_bld.zero,
+                                  &ctx->bld_base.uint_bld.zero,
                                   1, 0xf);
-               txq_emit(NULL, &ctx->soa.bld_base, &txq_emit_data);
+               txq_emit(NULL, &ctx->bld_base, &txq_emit_data);
 
                /* Compute -0.5 / size. */
                for (c = 0; c < 2; c++) {
@@ -4832,7 +4629,7 @@ static void si_lower_gather4_integer(struct si_shader_context *ctx,
                                                        LLVMConstInt(ctx->i32, c, 0), "");
                        half_texel[c] = LLVMBuildUIToFP(builder, half_texel[c], ctx->f32, "");
                        half_texel[c] =
-                               lp_build_emit_llvm_unary(&ctx->soa.bld_base,
+                               lp_build_emit_llvm_unary(&ctx->bld_base,
                                                         TGSI_OPCODE_RCP, half_texel[c]);
                        half_texel[c] = LLVMBuildFMul(builder, half_texel[c],
                                                      LLVMConstReal(ctx->f32, -0.5), "");
@@ -4987,35 +4784,6 @@ static void si_llvm_emit_txqs(
        emit_data->output[emit_data->chan] = samples;
 }
 
-/*
- * SI implements derivatives using the local data store (LDS)
- * All writes to the LDS happen in all executing threads at
- * the same time. TID is the Thread ID for the current
- * thread and is a value between 0 and 63, representing
- * the thread's position in the wavefront.
- *
- * For the pixel shader threads are grouped into quads of four pixels.
- * The TIDs of the pixels of a quad are:
- *
- *  +------+------+
- *  |4n + 0|4n + 1|
- *  +------+------+
- *  |4n + 2|4n + 3|
- *  +------+------+
- *
- * So, masking the TID with 0xfffffffc yields the TID of the top left pixel
- * of the quad, masking with 0xfffffffd yields the TID of the top pixel of
- * the current pixel's column, and masking with 0xfffffffe yields the TID
- * of the left pixel of the current pixel's row.
- *
- * Adding 1 yields the TID of the pixel to the right of the left pixel, and
- * adding 2 yields the TID of the pixel below the top pixel.
- */
-/* masks for thread ID. */
-#define TID_MASK_TOP_LEFT 0xfffffffc
-#define TID_MASK_TOP      0xfffffffd
-#define TID_MASK_LEFT     0xfffffffe
-
 static void si_llvm_emit_ddxy(
        const struct lp_build_tgsi_action *action,
        struct lp_build_tgsi_context *bld_base,
@@ -5024,59 +4792,24 @@ static void si_llvm_emit_ddxy(
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct gallivm_state *gallivm = bld_base->base.gallivm;
        unsigned opcode = emit_data->info->opcode;
-       LLVMValueRef thread_id, tl, trbl, tl_tid, trbl_tid, val, args[2];
+       LLVMValueRef val;
        int idx;
        unsigned mask;
 
-       thread_id = get_thread_id(ctx);
-
        if (opcode == TGSI_OPCODE_DDX_FINE)
-               mask = TID_MASK_LEFT;
+               mask = AC_TID_MASK_LEFT;
        else if (opcode == TGSI_OPCODE_DDY_FINE)
-               mask = TID_MASK_TOP;
+               mask = AC_TID_MASK_TOP;
        else
-               mask = TID_MASK_TOP_LEFT;
-
-       tl_tid = LLVMBuildAnd(gallivm->builder, thread_id,
-                               lp_build_const_int32(gallivm, mask), "");
+               mask = AC_TID_MASK_TOP_LEFT;
 
        /* for DDX we want to next X pixel, DDY next Y pixel. */
        idx = (opcode == TGSI_OPCODE_DDX || opcode == TGSI_OPCODE_DDX_FINE) ? 1 : 2;
-       trbl_tid = LLVMBuildAdd(gallivm->builder, tl_tid,
-                                 lp_build_const_int32(gallivm, idx), "");
 
        val = LLVMBuildBitCast(gallivm->builder, emit_data->args[0], ctx->i32, "");
-
-       if (ctx->screen->has_ds_bpermute) {
-               args[0] = LLVMBuildMul(gallivm->builder, tl_tid,
-                                      lp_build_const_int32(gallivm, 4), "");
-               args[1] = val;
-               tl = lp_build_intrinsic(gallivm->builder,
-                                       "llvm.amdgcn.ds.bpermute", ctx->i32,
-                                       args, 2, LP_FUNC_ATTR_READNONE);
-
-               args[0] = LLVMBuildMul(gallivm->builder, trbl_tid,
-                                      lp_build_const_int32(gallivm, 4), "");
-               trbl = lp_build_intrinsic(gallivm->builder,
-                                         "llvm.amdgcn.ds.bpermute", ctx->i32,
-                                         args, 2, LP_FUNC_ATTR_READNONE);
-       } else {
-               LLVMValueRef store_ptr, load_ptr0, load_ptr1;
-
-               store_ptr = build_gep0(ctx, ctx->lds, thread_id);
-               load_ptr0 = build_gep0(ctx, ctx->lds, tl_tid);
-               load_ptr1 = build_gep0(ctx, ctx->lds, trbl_tid);
-
-               LLVMBuildStore(gallivm->builder, val, store_ptr);
-               tl = LLVMBuildLoad(gallivm->builder, load_ptr0, "");
-               trbl = LLVMBuildLoad(gallivm->builder, load_ptr1, "");
-       }
-
-       tl = LLVMBuildBitCast(gallivm->builder, tl, ctx->f32, "");
-       trbl = LLVMBuildBitCast(gallivm->builder, trbl, ctx->f32, "");
-
-       emit_data->output[emit_data->chan] =
-               LLVMBuildFSub(gallivm->builder, trbl, tl, "");
+       val = ac_emit_ddxy(&ctx->ac, ctx->screen->has_ds_bpermute,
+                          mask, idx, ctx->lds, val);
+       emit_data->output[emit_data->chan] = val;
 }
 
 /*
@@ -5236,11 +4969,11 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
                                gallivm->builder, interp_param, uint->zero, "");
                        LLVMValueRef j = LLVMBuildExtractElement(
                                gallivm->builder, interp_param, uint->one, "");
-                       emit_data->output[chan] = build_fs_interp(bld_base,
+                       emit_data->output[chan] = ac_build_fs_interp(&ctx->ac,
                                llvm_chan, attr_number, params,
                                i, j);
                } else {
-                       emit_data->output[chan] = build_fs_interp_mov(bld_base,
+                       emit_data->output[chan] = ac_build_fs_interp_mov(&ctx->ac,
                                lp_build_const_int32(gallivm, 2), /* P0 */
                                llvm_chan, attr_number, params);
                }
@@ -5250,13 +4983,15 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
 static unsigned si_llvm_get_stream(struct lp_build_tgsi_context *bld_base,
                                       struct lp_build_emit_data *emit_data)
 {
-       LLVMValueRef (*imms)[4] = lp_soa_context(bld_base)->immediates;
+       struct si_shader_context *ctx = si_shader_context(bld_base);
        struct tgsi_src_register src0 = emit_data->inst->Src[0].Register;
+       LLVMValueRef imm;
        unsigned stream;
 
        assert(src0.File == TGSI_FILE_IMMEDIATE);
 
-       stream = LLVMConstIntGetZExtValue(imms[src0.Index][src0.SwizzleX]) & 0x3;
+       imm = ctx->imms[src0.Index * TGSI_NUM_CHANNELS + src0.SwizzleX];
+       stream = LLVMConstIntGetZExtValue(imm) & 0x3;
        return stream;
 }
 
@@ -5276,8 +5011,7 @@ static void si_llvm_emit_vertex(
                                            SI_PARAM_GS2VS_OFFSET);
        LLVMValueRef gs_next_vertex;
        LLVMValueRef can_emit, kill;
-       LLVMValueRef args[2];
-       unsigned chan;
+       unsigned chan, offset;
        int i;
        unsigned stream;
 
@@ -5312,28 +5046,33 @@ static void si_llvm_emit_vertex(
                lp_build_if(&if_state, gallivm, can_emit);
        }
 
+       offset = 0;
        for (i = 0; i < info->num_outputs; i++) {
-               LLVMValueRef *out_ptr =
-                       ctx->soa.outputs[i];
+               LLVMValueRef *out_ptr = ctx->outputs[i];
 
                for (chan = 0; chan < 4; chan++) {
+                       if (!(info->output_usagemask[i] & (1 << chan)) ||
+                           ((info->output_streams[i] >> (2 * chan)) & 3) != stream)
+                               continue;
+
                        LLVMValueRef out_val = LLVMBuildLoad(gallivm->builder, out_ptr[chan], "");
                        LLVMValueRef voffset =
-                               lp_build_const_int32(gallivm, (i * 4 + chan) *
+                               lp_build_const_int32(gallivm, offset *
                                                     shader->selector->gs_max_out_vertices);
+                       offset++;
 
                        voffset = lp_build_add(uint, voffset, gs_next_vertex);
                        voffset = lp_build_mul_imm(uint, voffset, 4);
 
                        out_val = LLVMBuildBitCast(gallivm->builder, out_val, ctx->i32, "");
 
-                       build_tbuffer_store(ctx,
-                                           ctx->gsvs_ring[stream],
-                                           out_val, 1,
-                                           voffset, soffset, 0,
-                                           V_008F0C_BUF_DATA_FORMAT_32,
-                                           V_008F0C_BUF_NUM_FORMAT_UINT,
-                                           1, 0, 1, 1, 0);
+                       ac_build_tbuffer_store(&ctx->ac,
+                                              ctx->gsvs_ring[stream],
+                                              out_val, 1,
+                                              voffset, soffset, 0,
+                                              V_008F0C_BUF_DATA_FORMAT_32,
+                                              V_008F0C_BUF_NUM_FORMAT_UINT,
+                                              1, 0, 1, 1, 0);
                }
        }
 
@@ -5343,11 +5082,8 @@ static void si_llvm_emit_vertex(
        LLVMBuildStore(gallivm->builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
 
        /* Signal vertex emission */
-       args[0] = lp_build_const_int32(gallivm, SENDMSG_GS_OP_EMIT | SENDMSG_GS | (stream << 8));
-       args[1] = LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID);
-       lp_build_intrinsic(gallivm->builder, "llvm.SI.sendmsg",
-                          ctx->voidt, args, 2, 0);
-
+       ac_emit_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
+                       LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));       
        if (!use_kill)
                lp_build_endif(&if_state);
 }
@@ -5359,16 +5095,12 @@ static void si_llvm_emit_primitive(
        struct lp_build_emit_data *emit_data)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef args[2];
        unsigned stream;
 
        /* Signal primitive cut */
        stream = si_llvm_get_stream(bld_base, emit_data);
-       args[0] = lp_build_const_int32(gallivm, SENDMSG_GS_OP_CUT | SENDMSG_GS | (stream << 8));
-       args[1] = LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID);
-       lp_build_intrinsic(gallivm->builder, "llvm.SI.sendmsg",
-                          ctx->voidt, args, 2, 0);
+       ac_emit_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8),
+                       LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
 }
 
 static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
@@ -5378,10 +5110,13 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct gallivm_state *gallivm = bld_base->base.gallivm;
 
-       /* The real barrier instruction isn’t needed, because an entire patch
+       /* SI only (thanks to a hw bug workaround):
+        * The real barrier instruction isn’t needed, because an entire patch
         * always fits into a single wave.
         */
-       if (ctx->type == PIPE_SHADER_TESS_CTRL) {
+       if (HAVE_LLVM >= 0x0309 &&
+           ctx->screen->b.chip_class == SI &&
+           ctx->type == PIPE_SHADER_TESS_CTRL) {
                emit_waitcnt(ctx, LGKM_CNT & VM_CNT);
                return;
        }
@@ -5449,20 +5184,6 @@ static void si_create_function(struct si_shader_context *ctx,
        }
 }
 
-static void create_meta_data(struct si_shader_context *ctx)
-{
-       struct gallivm_state *gallivm = ctx->soa.bld_base.base.gallivm;
-
-       ctx->invariant_load_md_kind = LLVMGetMDKindIDInContext(gallivm->context,
-                                                              "invariant.load", 14);
-       ctx->range_md_kind = LLVMGetMDKindIDInContext(gallivm->context,
-                                                    "range", 5);
-       ctx->uniform_md_kind = LLVMGetMDKindIDInContext(gallivm->context,
-                                                       "amdgpu.uniform", 14);
-
-       ctx->empty_md = LLVMMDNodeInContext(gallivm->context, NULL, 0);
-}
-
 static void declare_streamout_params(struct si_shader_context *ctx,
                                     struct pipe_stream_output_info *so,
                                     LLVMTypeRef *params, LLVMTypeRef i32,
@@ -5475,7 +5196,7 @@ static void declare_streamout_params(struct si_shader_context *ctx,
                if (ctx->type != PIPE_SHADER_TESS_EVAL)
                        params[ctx->param_streamout_config = (*num_params)++] = i32;
                else
-                       ctx->param_streamout_config = ctx->param_tess_offchip;
+                       ctx->param_streamout_config = *num_params - 1;
 
                params[ctx->param_streamout_write_index = (*num_params)++] = i32;
        }
@@ -5514,7 +5235,7 @@ static unsigned llvm_get_type_size(LLVMTypeRef type)
 static void declare_tess_lds(struct si_shader_context *ctx)
 {
        struct gallivm_state *gallivm = &ctx->gallivm;
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
        struct lp_build_context *uint = &bld_base->uint_bld;
 
        unsigned lds_size = ctx->screen->b.chip_class >= CIK ? 65536 : 32768;
@@ -5542,7 +5263,7 @@ static unsigned si_get_max_workgroup_size(struct si_shader *shader)
 
 static void create_function(struct si_shader_context *ctx)
 {
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
        struct gallivm_state *gallivm = bld_base->base.gallivm;
        struct si_shader *shader = ctx->shader;
        LLVMTypeRef params[SI_NUM_PARAMS + SI_NUM_VERTEX_BUFFERS], v3i32;
@@ -5639,10 +5360,10 @@ static void create_function(struct si_shader_context *ctx)
 
                if (shader->key.as_es) {
                        params[ctx->param_oc_lds = num_params++] = ctx->i32;
-                       params[ctx->param_tess_offchip = num_params++] = ctx->i32;
+                       params[num_params++] = ctx->i32;
                        params[ctx->param_es2gs_offset = num_params++] = ctx->i32;
                } else {
-                       params[ctx->param_tess_offchip = num_params++] = ctx->i32;
+                       params[num_params++] = ctx->i32;
                        declare_streamout_params(ctx, &shader->selector->so,
                                                 params, ctx->i32, &num_params);
                        params[ctx->param_oc_lds = num_params++] = ctx->i32;
@@ -5798,8 +5519,7 @@ static void create_function(struct si_shader_context *ctx)
                                                    LOCAL_ADDR_SPACE);
 
        if ((ctx->type == PIPE_SHADER_VERTEX && shader->key.as_ls) ||
-           ctx->type == PIPE_SHADER_TESS_CTRL ||
-           ctx->type == PIPE_SHADER_TESS_EVAL)
+           ctx->type == PIPE_SHADER_TESS_CTRL)
                declare_tess_lds(ctx);
 }
 
@@ -5809,8 +5529,8 @@ static void create_function(struct si_shader_context *ctx)
  */
 static void preload_ring_buffers(struct si_shader_context *ctx)
 {
-       struct gallivm_state *gallivm =
-               ctx->soa.bld_base.base.gallivm;
+       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+       LLVMBuilderRef builder = gallivm->builder;
 
        LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
                                            SI_PARAM_RW_BUFFERS);
@@ -5826,22 +5546,84 @@ static void preload_ring_buffers(struct si_shader_context *ctx)
                LLVMValueRef offset = lp_build_const_int32(gallivm, ring);
 
                ctx->esgs_ring =
-                       build_indexed_load_const(ctx, buf_ptr, offset);
+                       ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
        }
 
        if (ctx->shader->is_gs_copy_shader) {
-               LLVMValueRef offset = lp_build_const_int32(gallivm, SI_VS_RING_GSVS);
+               LLVMValueRef offset = lp_build_const_int32(gallivm, SI_RING_GSVS);
 
                ctx->gsvs_ring[0] =
-                       build_indexed_load_const(ctx, buf_ptr, offset);
-       }
-       if (ctx->type == PIPE_SHADER_GEOMETRY) {
-               int i;
-               for (i = 0; i < 4; i++) {
-                       LLVMValueRef offset = lp_build_const_int32(gallivm, SI_GS_RING_GSVS0 + i);
+                       ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
+       } else if (ctx->type == PIPE_SHADER_GEOMETRY) {
+               const struct si_shader_selector *sel = ctx->shader->selector;
+               struct lp_build_context *uint = &ctx->bld_base.uint_bld;
+               LLVMValueRef offset = lp_build_const_int32(gallivm, SI_RING_GSVS);
+               LLVMValueRef base_ring;
+
+               base_ring = ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
+
+               /* The conceptual layout of the GSVS ring is
+                *   v0c0 .. vLv0 v0c1 .. vLc1 ..
+                * but the real memory layout is swizzled across
+                * threads:
+                *   t0v0c0 .. t15v0c0 t0v1c0 .. t15v1c0 ... t15vLcL
+                *   t16v0c0 ..
+                * Override the buffer descriptor accordingly.
+                */
+               LLVMTypeRef v2i64 = LLVMVectorType(ctx->i64, 2);
+               uint64_t stream_offset = 0;
+
+               for (unsigned stream = 0; stream < 4; ++stream) {
+                       unsigned num_components;
+                       unsigned stride;
+                       unsigned num_records;
+                       LLVMValueRef ring, tmp;
 
-                       ctx->gsvs_ring[i] =
-                               build_indexed_load_const(ctx, buf_ptr, offset);
+                       num_components = sel->info.num_stream_output_components[stream];
+                       if (!num_components)
+                               continue;
+
+                       stride = 4 * num_components * sel->gs_max_out_vertices;
+
+                       /* Limit on the stride field for <= CIK. */
+                       assert(stride < (1 << 14));
+
+                       num_records = 64;
+
+                       ring = LLVMBuildBitCast(builder, base_ring, v2i64, "");
+                       tmp = LLVMBuildExtractElement(builder, ring, uint->zero, "");
+                       tmp = LLVMBuildAdd(builder, tmp,
+                                          LLVMConstInt(ctx->i64,
+                                                       stream_offset, 0), "");
+                       stream_offset += stride * 64;
+
+                       ring = LLVMBuildInsertElement(builder, ring, tmp, uint->zero, "");
+                       ring = LLVMBuildBitCast(builder, ring, ctx->v4i32, "");
+                       tmp = LLVMBuildExtractElement(builder, ring, uint->one, "");
+                       tmp = LLVMBuildOr(builder, tmp,
+                               LLVMConstInt(ctx->i32,
+                                            S_008F04_STRIDE(stride) |
+                                            S_008F04_SWIZZLE_ENABLE(1), 0), "");
+                       ring = LLVMBuildInsertElement(builder, ring, tmp, uint->one, "");
+                       ring = LLVMBuildInsertElement(builder, ring,
+                                       LLVMConstInt(ctx->i32, num_records, 0),
+                                       LLVMConstInt(ctx->i32, 2, 0), "");
+                       ring = LLVMBuildInsertElement(builder, ring,
+                               LLVMConstInt(ctx->i32,
+                                            S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) |
+                                            S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) |
+                                            S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) |
+                                            S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W) |
+                                            S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) |
+                                            S_008F0C_DATA_FORMAT(V_008F0C_BUF_DATA_FORMAT_32) |
+                                            S_008F0C_ELEMENT_SIZE(1) | /* element_size = 4 (bytes) */
+                                            S_008F0C_INDEX_STRIDE(1) | /* index_stride = 16 (elements) */
+                                            S_008F0C_ADD_TID_ENABLE(1),
+                                            0),
+                               LLVMConstInt(ctx->i32, 3, 0), "");
+                       ring = LLVMBuildBitCast(builder, ring, ctx->v16i8, "");
+
+                       ctx->gsvs_ring[stream] = ring;
                }
        }
 }
@@ -5850,8 +5632,7 @@ static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx,
                                         LLVMValueRef param_rw_buffers,
                                         unsigned param_pos_fixed_pt)
 {
-       struct lp_build_tgsi_context *bld_base =
-               &ctx->soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
        struct gallivm_state *gallivm = bld_base->base.gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        LLVMValueRef slot, desc, offset, row, bit, address[2];
@@ -5865,7 +5646,7 @@ static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx,
 
        /* Load the buffer descriptor. */
        slot = lp_build_const_int32(gallivm, SI_PS_CONST_POLY_STIPPLE);
-       desc = build_indexed_load_const(ctx, param_rw_buffers, slot);
+       desc = ac_build_indexed_load_const(&ctx->ac, param_rw_buffers, slot);
 
        /* The stipple pattern is 32x32, each row has 32 bits. */
        offset = LLVMBuildMul(builder, address[1],
@@ -6025,7 +5806,8 @@ int si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader)
        r600_resource_reference(&shader->bo, NULL);
        shader->bo = (struct r600_resource*)
                     pipe_buffer_create(&sscreen->b.b, 0,
-                                       PIPE_USAGE_IMMUTABLE, bo_size);
+                                       PIPE_USAGE_IMMUTABLE,
+                                       align(bo_size, SI_CPDMA_ALIGNMENT));
        if (!shader->bo)
                return -ENOMEM;
 
@@ -6102,7 +5884,8 @@ static void si_shader_dump_stats(struct si_screen *sscreen,
                                 struct si_shader *shader,
                                 struct pipe_debug_callback *debug,
                                 unsigned processor,
-                                FILE *file)
+                                FILE *file,
+                                bool check_debug_option)
 {
        struct si_shader_config *conf = &shader->config;
        unsigned num_inputs = shader->selector ? shader->selector->info.num_inputs : 0;
@@ -6153,7 +5936,7 @@ static void si_shader_dump_stats(struct si_screen *sscreen,
        if (lds_per_wave)
                max_simd_waves = MIN2(max_simd_waves, 16384 / lds_per_wave);
 
-       if (file != stderr ||
+       if (!check_debug_option ||
            r600_can_dump_shader(&sscreen->b, processor)) {
                if (processor == PIPE_SHADER_FRAGMENT) {
                        fprintf(file, "*** SHADER CONFIG ***\n"
@@ -6190,8 +5973,7 @@ static void si_shader_dump_stats(struct si_screen *sscreen,
                           conf->spilled_vgprs, conf->private_mem_vgprs);
 }
 
-static const char *si_get_shader_name(struct si_shader *shader,
-                                     unsigned processor)
+const char *si_get_shader_name(struct si_shader *shader, unsigned processor)
 {
        switch (processor) {
        case PIPE_SHADER_VERTEX:
@@ -6224,19 +6006,19 @@ static const char *si_get_shader_name(struct si_shader *shader,
 
 void si_shader_dump(struct si_screen *sscreen, struct si_shader *shader,
                    struct pipe_debug_callback *debug, unsigned processor,
-                   FILE *file)
+                   FILE *file, bool check_debug_option)
 {
-       if (file != stderr ||
+       if (!check_debug_option ||
            r600_can_dump_shader(&sscreen->b, processor))
                si_dump_shader_key(processor, &shader->key, file);
 
-       if (file != stderr && shader->binary.llvm_ir_string) {
+       if (!check_debug_option && shader->binary.llvm_ir_string) {
                fprintf(file, "\n%s - main shader part - LLVM IR:\n\n",
                        si_get_shader_name(shader, processor));
                fprintf(file, "%s\n", shader->binary.llvm_ir_string);
        }
 
-       if (file != stderr ||
+       if (!check_debug_option ||
            (r600_can_dump_shader(&sscreen->b, processor) &&
             !(sscreen->b.debug_flags & DBG_NO_ASM))) {
                fprintf(file, "\n%s:\n", si_get_shader_name(shader, processor));
@@ -6253,7 +6035,8 @@ void si_shader_dump(struct si_screen *sscreen, struct si_shader *shader,
                fprintf(file, "\n");
        }
 
-       si_shader_dump_stats(sscreen, shader, debug, processor, file);
+       si_shader_dump_stats(sscreen, shader, debug, processor, file,
+                            check_debug_option);
 }
 
 int si_compile_llvm(struct si_screen *sscreen,
@@ -6273,7 +6056,7 @@ int si_compile_llvm(struct si_screen *sscreen,
 
                if (!(sscreen->b.debug_flags & (DBG_NO_IR | DBG_PREOPT_IR))) {
                        fprintf(stderr, "%s LLVM IR:\n\n", name);
-                       LLVMDumpModule(mod);
+                       ac_dump_module(mod);
                        fprintf(stderr, "\n");
                }
        }
@@ -6345,7 +6128,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
        struct si_shader *shader;
        struct gallivm_state *gallivm = &ctx.gallivm;
        LLVMBuilderRef builder;
-       struct lp_build_tgsi_context *bld_base = &ctx.soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx.bld_base;
        struct lp_build_context *uint = &bld_base->uint_bld;
        struct si_shader_output_values *outputs;
        struct tgsi_shader_info *gsinfo = &gs_selector->info;
@@ -6372,7 +6155,6 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
 
        builder = gallivm->builder;
 
-       create_meta_data(&ctx);
        create_function(&ctx);
        preload_ring_buffers(&ctx);
 
@@ -6396,33 +6178,26 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
        else
                stream_id = uint->zero;
 
-       /* Fetch vertex data from GSVS ring */
+       /* Fill in output information. */
        for (i = 0; i < gsinfo->num_outputs; ++i) {
-               unsigned chan;
-
                outputs[i].semantic_name = gsinfo->output_semantic_name[i];
                outputs[i].semantic_index = gsinfo->output_semantic_index[i];
 
-               for (chan = 0; chan < 4; chan++) {
+               for (int chan = 0; chan < 4; chan++) {
                        outputs[i].vertex_stream[chan] =
                                (gsinfo->output_streams[i] >> (2 * chan)) & 3;
-
-                       args[2] = lp_build_const_int32(gallivm,
-                                                      (i * 4 + chan) *
-                                                      gs_selector->gs_max_out_vertices * 16 * 4);
-
-                       outputs[i].values[chan] =
-                               LLVMBuildBitCast(gallivm->builder,
-                                                lp_build_intrinsic(gallivm->builder,
-                                                                "llvm.SI.buffer.load.dword.i32.i32",
-                                                                ctx.i32, args, 9,
-                                                                LP_FUNC_ATTR_READONLY),
-                                                ctx.f32, "");
                }
        }
 
+       LLVMBasicBlockRef end_bb;
+       LLVMValueRef switch_inst;
+
+       end_bb = LLVMAppendBasicBlockInContext(gallivm->context, ctx.main_fn, "end");
+       switch_inst = LLVMBuildSwitch(builder, stream_id, end_bb, 4);
+
        for (int stream = 0; stream < 4; stream++) {
-               struct lp_build_if_state if_ctx_stream;
+               LLVMBasicBlockRef bb;
+               unsigned offset;
 
                if (!gsinfo->num_stream_output_components[stream])
                        continue;
@@ -6430,13 +6205,36 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
                if (stream > 0 && !gs_selector->so.num_outputs)
                        continue;
 
-               LLVMValueRef is_stream =
-                       LLVMBuildICmp(builder, LLVMIntEQ,
-                                     stream_id,
-                                     lp_build_const_int32(gallivm, stream), "");
+               bb = LLVMInsertBasicBlockInContext(gallivm->context, end_bb, "out");
+               LLVMAddCase(switch_inst, lp_build_const_int32(gallivm, stream), bb);
+               LLVMPositionBuilderAtEnd(builder, bb);
+
+               /* Fetch vertex data from GSVS ring */
+               offset = 0;
+               for (i = 0; i < gsinfo->num_outputs; ++i) {
+                       for (unsigned chan = 0; chan < 4; chan++) {
+                               if (!(gsinfo->output_usagemask[i] & (1 << chan)) ||
+                                   outputs[i].vertex_stream[chan] != stream) {
+                                       outputs[i].values[chan] = ctx.bld_base.base.undef;
+                                       continue;
+                               }
 
-               lp_build_if(&if_ctx_stream, gallivm, is_stream);
+                               args[2] = lp_build_const_int32(
+                                       gallivm,
+                                       offset * gs_selector->gs_max_out_vertices * 16 * 4);
+                               offset++;
 
+                               outputs[i].values[chan] =
+                                       LLVMBuildBitCast(gallivm->builder,
+                                                lp_build_intrinsic(gallivm->builder,
+                                                                "llvm.SI.buffer.load.dword.i32.i32",
+                                                                ctx.i32, args, 9,
+                                                                LP_FUNC_ATTR_READONLY),
+                                                ctx.f32, "");
+                       }
+               }
+
+               /* Streamout and exports. */
                if (gs_selector->so.num_outputs) {
                        si_llvm_emit_streamout(&ctx, outputs,
                                               gsinfo->num_outputs,
@@ -6446,15 +6244,17 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
                if (stream == 0)
                        si_llvm_export_vs(bld_base, outputs, gsinfo->num_outputs);
 
-               lp_build_endif(&if_ctx_stream);
+               LLVMBuildBr(builder, end_bb);
        }
 
+       LLVMPositionBuilderAtEnd(builder, end_bb);
+
        LLVMBuildRetVoid(gallivm->builder);
 
        /* Dump LLVM IR before any optimization passes */
        if (sscreen->b.debug_flags & DBG_PREOPT_IR &&
            r600_can_dump_shader(&sscreen->b, PIPE_SHADER_GEOMETRY))
-               LLVMDumpModule(bld_base->base.gallivm->module);
+               ac_dump_module(bld_base->base.gallivm->module);
 
        si_llvm_finalize_module(&ctx,
                r600_extra_shader_checks(&sscreen->b, PIPE_SHADER_GEOMETRY));
@@ -6468,7 +6268,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
                if (r600_can_dump_shader(&sscreen->b, PIPE_SHADER_GEOMETRY))
                        fprintf(stderr, "GS Copy Shader:\n");
                si_shader_dump(sscreen, ctx.shader, debug,
-                              PIPE_SHADER_GEOMETRY, stderr);
+                              PIPE_SHADER_GEOMETRY, stderr, true);
                r = si_shader_binary_upload(sscreen, ctx.shader);
        }
 
@@ -6500,7 +6300,11 @@ static void si_dump_shader_key(unsigned shader, struct si_shader_key *key,
                fprintf(f, "  part.vs.epilog.export_prim_id = %u\n", key->part.vs.epilog.export_prim_id);
                fprintf(f, "  as_es = %u\n", key->as_es);
                fprintf(f, "  as_ls = %u\n", key->as_ls);
-               fprintf(f, "  mono.vs.fix_fetch = 0x%x\n", key->mono.vs.fix_fetch);
+
+               fprintf(f, "  mono.vs.fix_fetch = {");
+               for (i = 0; i < SI_MAX_ATTRIBS; i++)
+                       fprintf(f, !i ? "%u" : ", %u", key->mono.vs.fix_fetch[i]);
+               fprintf(f, "}\n");
                break;
 
        case PIPE_SHADER_TESS_CTRL:
@@ -6565,7 +6369,7 @@ static void si_init_shader_ctx(struct si_shader_context *ctx,
                (shader && shader->selector) ? &shader->selector->info : NULL,
                (shader && shader->selector) ? shader->selector->tokens : NULL);
 
-       bld_base = &ctx->soa.bld_base;
+       bld_base = &ctx->bld_base;
        bld_base->emit_fetch_funcs[TGSI_FILE_CONSTANT] = fetch_constant;
 
        bld_base->op_actions[TGSI_OPCODE_INTERP_CENTROID] = interp_action;
@@ -6811,7 +6615,7 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx,
                                 struct si_shader *shader)
 {
        struct si_shader_selector *sel = shader->selector;
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
 
        switch (ctx->type) {
        case PIPE_SHADER_VERTEX:
@@ -6852,7 +6656,6 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx,
                return false;
        }
 
-       create_meta_data(ctx);
        create_function(ctx);
        preload_ring_buffers(ctx);
 
@@ -7360,7 +7163,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 
        shader->info.uses_instanceid = sel->info.uses_instanceid;
 
-       bld_base = &ctx.soa.bld_base;
+       bld_base = &ctx.bld_base;
        ctx.load_system_value = declare_system_value;
 
        if (!si_compile_tgsi_main(&ctx, shader)) {
@@ -7458,7 +7261,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
        /* Dump LLVM IR before any optimization passes */
        if (sscreen->b.debug_flags & DBG_PREOPT_IR &&
            r600_can_dump_shader(&sscreen->b, ctx.type))
-               LLVMDumpModule(mod);
+               ac_dump_module(mod);
 
        si_llvm_finalize_module(&ctx,
                                    r600_extra_shader_checks(&sscreen->b, ctx.type));
@@ -7759,7 +7562,7 @@ static void si_build_vs_epilog_function(struct si_shader_context *ctx,
                                        union si_shader_part_key *key)
 {
        struct gallivm_state *gallivm = &ctx->gallivm;
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
        LLVMTypeRef params[5];
        int num_params, i;
 
@@ -7880,7 +7683,7 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx,
                                         union si_shader_part_key *key)
 {
        struct gallivm_state *gallivm = &ctx->gallivm;
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
        LLVMTypeRef params[16];
        LLVMValueRef func;
        int last_sgpr, num_params;
@@ -8215,7 +8018,7 @@ static void si_build_ps_epilog_function(struct si_shader_context *ctx,
                                        union si_shader_part_key *key)
 {
        struct gallivm_state *gallivm = &ctx->gallivm;
-       struct lp_build_tgsi_context *bld_base = &ctx->soa.bld_base;
+       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
        LLVMTypeRef params[16+8*4+3];
        LLVMValueRef depth = NULL, stencil = NULL, samplemask = NULL;
        int last_sgpr, num_params, i;
@@ -8437,7 +8240,7 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
                     struct pipe_debug_callback *debug)
 {
        struct si_shader_selector *sel = shader->selector;
-       struct si_shader *mainp = sel->main_shader_part;
+       struct si_shader *mainp = *si_get_main_shader_part(sel, &shader->key);
        int r;
 
        /* LS, ES, VS are compiled on demand if the main part hasn't been
@@ -8526,7 +8329,7 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
 
        si_fix_resource_usage(sscreen, shader);
        si_shader_dump(sscreen, shader, debug, sel->info.processor,
-                      stderr);
+                      stderr, true);
 
        /* Upload. */
        r = si_shader_binary_upload(sscreen, shader);