radeonsi: replace si_shader_context::soa by bld_base
[mesa.git] / src / gallium / drivers / radeonsi / si_shader.c
index 8827643caef98ba0a050b905c5a01db0971a6781..d45c0e8649c87c90f387adbf97801c5da626a010 100644 (file)
@@ -55,8 +55,9 @@ static const char *scratch_rsrc_dword1_symbol =
 struct si_shader_output_values
 {
        LLVMValueRef values[4];
-       unsigned name;
-       unsigned sid;
+       unsigned semantic_name;
+       unsigned semantic_index;
+       ubyte vertex_stream[4];
 };
 
 static void si_init_shader_ctx(struct si_shader_context *ctx,
@@ -175,7 +176,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)
@@ -250,7 +251,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),
@@ -260,7 +261,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),
@@ -321,7 +322,7 @@ 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 lp_build_tgsi_context *bld_base = &ctx->bld_base;
        struct gallivm_state *gallivm = bld_base->base.gallivm;
 
        LLVMBuildStore(gallivm->builder, value,
@@ -341,7 +342,7 @@ 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 lp_build_tgsi_context *bld_base = &ctx->bld_base;
        struct gallivm_state *gallivm = bld_base->base.gallivm;
        LLVMValueRef pointer;
 
@@ -369,8 +370,8 @@ static LLVMValueRef get_instance_index_for_fetch(
        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);
@@ -390,7 +391,7 @@ static void declare_input_vs(
        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;
@@ -507,10 +508,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), "");
@@ -547,7 +548,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;
@@ -645,7 +646,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;
 
@@ -689,7 +690,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;
@@ -1121,7 +1122,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];
@@ -1319,7 +1320,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;
@@ -1432,9 +1433,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;
@@ -1478,7 +1479,7 @@ 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);
 }
 
@@ -1543,8 +1544,8 @@ 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);
@@ -1571,8 +1572,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;
 
@@ -1626,7 +1627,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)),
                };
@@ -1649,9 +1650,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;
@@ -1707,7 +1708,7 @@ static void declare_system_value(
                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;
@@ -1732,7 +1733,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:
@@ -1770,16 +1771,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");
@@ -1793,7 +1797,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;
 
@@ -1859,7 +1863,7 @@ static LLVMValueRef fetch_constant(
                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,
@@ -1914,8 +1918,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;
@@ -2150,7 +2153,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;
@@ -2213,31 +2216,72 @@ static void si_dump_streamout(struct pipe_stream_output_info *so)
        }
 }
 
-/* On SI, the vertex shader is responsible for writing streamout data
- * to buffers. */
-static void si_llvm_emit_streamout(struct si_shader_context *ctx,
-                                  struct si_shader_output_values *outputs,
-                                  unsigned noutput)
+static void emit_streamout_output(struct si_shader_context *ctx,
+                                 LLVMValueRef const *so_buffers,
+                                 LLVMValueRef const *so_write_offsets,
+                                 struct pipe_stream_output *stream_out,
+                                 struct si_shader_output_values *shader_out)
 {
-       struct pipe_stream_output_info *so = &ctx->shader->selector->so;
        struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
-       int i, j;
-       struct lp_build_if_state if_ctx;
-       LLVMValueRef so_buffers[4];
-       LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
-                                           SI_PARAM_RW_BUFFERS);
+       unsigned buf_idx = stream_out->output_buffer;
+       unsigned start = stream_out->start_component;
+       unsigned num_comps = stream_out->num_components;
+       LLVMValueRef out[4];
 
-       /* Load the descriptors. */
-       for (i = 0; i < 4; ++i) {
-               if (ctx->shader->selector->so.stride[i]) {
-                       LLVMValueRef offset = lp_build_const_int32(gallivm,
-                                                                  SI_VS_STREAMOUT_BUF0 + i);
+       assert(num_comps && num_comps <= 4);
+       if (!num_comps || num_comps > 4)
+               return;
 
-                       so_buffers[i] = build_indexed_load_const(ctx, buf_ptr, offset);
+       /* Load the output as int. */
+       for (int j = 0; j < num_comps; j++) {
+               assert(stream_out->stream == shader_out->vertex_stream[start + j]);
+
+               out[j] = LLVMBuildBitCast(builder,
+                                         shader_out->values[start + j],
+                               ctx->i32, "");
+       }
+
+       /* Pack the output. */
+       LLVMValueRef vdata = NULL;
+
+       switch (num_comps) {
+       case 1: /* as i32 */
+               vdata = out[0];
+               break;
+       case 2: /* as v2i32 */
+       case 3: /* as v4i32 (aligned to 4) */
+       case 4: /* as v4i32 */
+               vdata = LLVMGetUndef(LLVMVectorType(ctx->i32, util_next_power_of_two(num_comps)));
+               for (int j = 0; j < num_comps; j++) {
+                       vdata = LLVMBuildInsertElement(builder, vdata, out[j],
+                                                      LLVMConstInt(ctx->i32, j, 0), "");
                }
+               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);
+}
+
+/**
+ * Write streamout data to buffers for vertex stream @p stream (different
+ * vertex streams can occur for GS copy shaders).
+ */
+static void si_llvm_emit_streamout(struct si_shader_context *ctx,
+                                  struct si_shader_output_values *outputs,
+                                  unsigned noutput, unsigned stream)
+{
+       struct si_shader_selector *sel = ctx->shader->selector;
+       struct pipe_stream_output_info *so = &sel->so;
+       struct gallivm_state *gallivm = &ctx->gallivm;
+       LLVMBuilderRef builder = gallivm->builder;
+       int i;
+       struct lp_build_if_state if_ctx;
+
        /* Get bits [22:16], i.e. (so_param >> 16) & 127; */
        LLVMValueRef so_vtx_count =
                unpack_param(ctx, ctx->param_streamout_config, 16, 7);
@@ -2248,9 +2292,6 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
        LLVMValueRef can_emit =
                LLVMBuildICmp(builder, LLVMIntULT, tid, so_vtx_count, "");
 
-       LLVMValueRef stream_id =
-               unpack_param(ctx, ctx->param_streamout_config, 24, 2);
-
        /* Emit the streamout code conditionally. This actually avoids
         * out-of-bounds buffer access. The hw tells us via the SGPR
         * (so_vtx_count) which threads are allowed to emit streamout data. */
@@ -2269,12 +2310,22 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
                /* Compute (streamout_write_index + thread_id). */
                so_write_index = LLVMBuildAdd(builder, so_write_index, tid, "");
 
-               /* Compute the write offset for each enabled buffer. */
+               /* Load the descriptor and compute the write offset for each
+                * enabled buffer. */
                LLVMValueRef so_write_offset[4] = {};
+               LLVMValueRef so_buffers[4];
+               LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
+                                                   SI_PARAM_RW_BUFFERS);
+
                for (i = 0; i < 4; i++) {
                        if (!so->stride[i])
                                continue;
 
+                       LLVMValueRef offset = lp_build_const_int32(gallivm,
+                                                                  SI_VS_STREAMOUT_BUF0 + i);
+
+                       so_buffers[i] = build_indexed_load_const(ctx, buf_ptr, offset);
+
                        LLVMValueRef so_offset = LLVMGetParam(ctx->main_fn,
                                                              ctx->param_streamout_offset[i]);
                        so_offset = LLVMBuildMul(builder, so_offset, LLVMConstInt(ctx->i32, 4, 0), "");
@@ -2286,58 +2337,16 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
 
                /* Write streamout data. */
                for (i = 0; i < so->num_outputs; i++) {
-                       unsigned buf_idx = so->output[i].output_buffer;
                        unsigned reg = so->output[i].register_index;
-                       unsigned start = so->output[i].start_component;
-                       unsigned num_comps = so->output[i].num_components;
-                       unsigned stream = so->output[i].stream;
-                       LLVMValueRef out[4];
-                       struct lp_build_if_state if_ctx_stream;
-
-                       assert(num_comps && num_comps <= 4);
-                       if (!num_comps || num_comps > 4)
-                               continue;
 
                        if (reg >= noutput)
                                continue;
 
-                       /* Load the output as int. */
-                       for (j = 0; j < num_comps; j++) {
-                               out[j] = LLVMBuildBitCast(builder,
-                                                         outputs[reg].values[start+j],
-                                               ctx->i32, "");
-                       }
-
-                       /* Pack the output. */
-                       LLVMValueRef vdata = NULL;
-
-                       switch (num_comps) {
-                       case 1: /* as i32 */
-                               vdata = out[0];
-                               break;
-                       case 2: /* as v2i32 */
-                       case 3: /* as v4i32 (aligned to 4) */
-                       case 4: /* as v4i32 */
-                               vdata = LLVMGetUndef(LLVMVectorType(ctx->i32, util_next_power_of_two(num_comps)));
-                               for (j = 0; j < num_comps; j++) {
-                                       vdata = LLVMBuildInsertElement(builder, vdata, out[j],
-                                                                      LLVMConstInt(ctx->i32, j, 0), "");
-                               }
-                               break;
-                       }
+                       if (stream != so->output[i].stream)
+                               continue;
 
-                       LLVMValueRef can_emit_stream =
-                               LLVMBuildICmp(builder, LLVMIntEQ,
-                                             stream_id,
-                                             lp_build_const_int32(gallivm, stream), "");
-
-                       lp_build_if(&if_ctx_stream, gallivm, can_emit_stream);
-                       build_tbuffer_store_dwords(ctx, so_buffers[buf_idx],
-                                                  vdata, num_comps,
-                                                  so_write_offset[buf_idx],
-                                                  LLVMConstInt(ctx->i32, 0, 0),
-                                                  so->output[i].dst_offset*4);
-                       lp_build_endif(&if_ctx_stream);
+                       emit_streamout_output(ctx, so_buffers, so_write_offset,
+                                             &so->output[i], &outputs[reg]);
                }
        }
        lp_build_endif(&if_ctx);
@@ -2352,8 +2361,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;
@@ -2363,13 +2371,9 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base,
        unsigned pos_idx;
        int i;
 
-       if (outputs && ctx->shader->selector->so.num_outputs) {
-               si_llvm_emit_streamout(ctx, outputs, noutput);
-       }
-
        for (i = 0; i < noutput; i++) {
-               semantic_name = outputs[i].name;
-               semantic_index = outputs[i].sid;
+               semantic_name = outputs[i].semantic_name;
+               semantic_index = outputs[i].semantic_index;
                bool export_param = true;
 
                switch (semantic_name) {
@@ -2391,6 +2395,12 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base,
                        break;
                }
 
+               if (outputs[i].vertex_stream[0] != 0 &&
+                   outputs[i].vertex_stream[1] != 0 &&
+                   outputs[i].vertex_stream[2] != 0 &&
+                   outputs[i].vertex_stream[3] != 0)
+                       export_param = false;
+
 handle_semantic:
                /* Select the correct target */
                switch(semantic_name) {
@@ -2411,15 +2421,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;
@@ -2432,6 +2433,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:
@@ -2774,7 +2777,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);
@@ -2800,8 +2803,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 ||
@@ -2878,7 +2880,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);
@@ -2890,14 +2892,18 @@ static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
        }
 
        for (i = 0; i < info->num_outputs; i++) {
-               outputs[i].name = info->output_semantic_name[i];
-               outputs[i].sid = info->output_semantic_index[i];
+               outputs[i].semantic_name = info->output_semantic_name[i];
+               outputs[i].semantic_index = info->output_semantic_index[i];
 
-               for (j = 0; j < 4; j++)
+               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;
+               }
+
        }
 
        /* Return the primitive ID from the LLVM function. */
@@ -2908,6 +2914,8 @@ static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
                                             get_primitive_id(bld_base, 0)),
                                     VS_EPILOG_PRIMID_LOC, "");
 
+       if (ctx->shader->selector->so.num_outputs)
+               si_llvm_emit_streamout(ctx, outputs, i, 0);
        si_llvm_export_vs(bld_base, outputs, i);
        FREE(outputs);
 }
@@ -3140,22 +3148,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",
@@ -3651,7 +3659,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;
@@ -3667,7 +3675,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;
@@ -3833,7 +3841,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;
@@ -3904,7 +3912,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;
@@ -3916,7 +3924,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);
@@ -4034,12 +4042,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, "");
@@ -4595,7 +4603,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++)
@@ -4709,7 +4721,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);
@@ -4717,7 +4728,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:
@@ -4727,7 +4738,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:
@@ -4735,7 +4746,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 */
                        }
@@ -4755,13 +4766,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);
                }
@@ -4806,9 +4816,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++) {
@@ -4817,7 +4827,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), "");
@@ -5235,13 +5245,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;
 }
 
@@ -5262,7 +5274,7 @@ static void si_llvm_emit_vertex(
        LLVMValueRef gs_next_vertex;
        LLVMValueRef can_emit, kill;
        LLVMValueRef args[2];
-       unsigned chan;
+       unsigned chan, offset;
        int i;
        unsigned stream;
 
@@ -5281,7 +5293,7 @@ static void si_llvm_emit_vertex(
         * further memory loads and may allow LLVM to skip to the end
         * altogether.
         */
-       can_emit = LLVMBuildICmp(gallivm->builder, LLVMIntULE, gs_next_vertex,
+       can_emit = LLVMBuildICmp(gallivm->builder, LLVMIntULT, gs_next_vertex,
                                 lp_build_const_int32(gallivm,
                                                      shader->selector->gs_max_out_vertices), "");
 
@@ -5297,15 +5309,20 @@ 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);
@@ -5436,7 +5453,7 @@ 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;
+       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
 
        ctx->invariant_load_md_kind = LLVMGetMDKindIDInContext(gallivm->context,
                                                               "invariant.load", 14);
@@ -5499,7 +5516,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;
@@ -5527,7 +5544,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;
@@ -5794,8 +5811,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);
@@ -5815,18 +5832,80 @@ static void preload_ring_buffers(struct si_shader_context *ctx)
        }
 
        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);
+       } 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 = build_indexed_load_const(ctx, 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;
                }
        }
 }
@@ -5835,8 +5914,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];
@@ -6087,7 +6165,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;
@@ -6138,7 +6217,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"
@@ -6209,19 +6288,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));
@@ -6238,7 +6317,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,
@@ -6329,7 +6409,8 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
        struct si_shader_context ctx;
        struct si_shader *shader;
        struct gallivm_state *gallivm = &ctx.gallivm;
-       struct lp_build_tgsi_context *bld_base = &ctx.soa.bld_base;
+       LLVMBuilderRef builder;
+       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;
@@ -6354,6 +6435,8 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
        si_init_shader_ctx(&ctx, sscreen, shader, tm);
        ctx.type = PIPE_SHADER_VERTEX;
 
+       builder = gallivm->builder;
+
        create_meta_data(&ctx);
        create_function(&ctx);
        preload_ring_buffers(&ctx);
@@ -6370,29 +6453,84 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
        args[7] = uint->one;  /* SLC */
        args[8] = uint->zero; /* TFE */
 
-       /* Fetch vertex data from GSVS ring */
+       /* Fetch the vertex stream ID.*/
+       LLVMValueRef stream_id;
+
+       if (gs_selector->so.num_outputs)
+               stream_id = unpack_param(&ctx, ctx.param_streamout_config, 24, 2);
+       else
+               stream_id = uint->zero;
+
+       /* 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];
 
-               outputs[i].name = gsinfo->output_semantic_name[i];
-               outputs[i].sid = gsinfo->output_semantic_index[i];
+               for (int chan = 0; chan < 4; chan++) {
+                       outputs[i].vertex_stream[chan] =
+                               (gsinfo->output_streams[i] >> (2 * chan)) & 3;
+               }
+       }
 
-               for (chan = 0; chan < 4; chan++) {
-                       args[2] = lp_build_const_int32(gallivm,
-                                                      (i * 4 + chan) *
-                                                      gs_selector->gs_max_out_vertices * 16 * 4);
+       LLVMBasicBlockRef end_bb;
+       LLVMValueRef switch_inst;
 
-                       outputs[i].values[chan] =
-                               LLVMBuildBitCast(gallivm->builder,
+       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++) {
+               LLVMBasicBlockRef bb;
+               unsigned offset;
+
+               if (!gsinfo->num_stream_output_components[stream])
+                       continue;
+
+               if (stream > 0 && !gs_selector->so.num_outputs)
+                       continue;
+
+               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;
+                               }
+
+                               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,
+                                              stream);
                }
+
+               if (stream == 0)
+                       si_llvm_export_vs(bld_base, outputs, gsinfo->num_outputs);
+
+               LLVMBuildBr(builder, end_bb);
        }
 
-       si_llvm_export_vs(bld_base, outputs, gsinfo->num_outputs);
+       LLVMPositionBuilderAtEnd(builder, end_bb);
 
        LLVMBuildRetVoid(gallivm->builder);
 
@@ -6413,7 +6551,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);
        }
 
@@ -6510,7 +6648,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;
@@ -6756,7 +6894,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:
@@ -7305,7 +7443,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)) {
@@ -7704,7 +7842,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;
 
@@ -7825,7 +7963,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;
@@ -8160,7 +8298,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;
@@ -8471,7 +8609,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);