radeonsi: change the bit-packing of LS out/TCS in data
[mesa.git] / src / gallium / drivers / radeonsi / si_shader.c
index 4a346538ee4049e4f9005c5a78f61a3f2769e936..5c17c640a3c9462500594b3f040eea3f6c376a9b 100644 (file)
@@ -72,6 +72,8 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
 static void si_dump_shader_key(unsigned shader, struct si_shader_key *key,
                               FILE *f);
 
+static unsigned llvm_get_type_size(LLVMTypeRef type);
+
 static void si_build_vs_prolog_function(struct si_shader_context *ctx,
                                        union si_shader_part_key *key);
 static void si_build_vs_epilog_function(struct si_shader_context *ctx,
@@ -173,12 +175,12 @@ static LLVMValueRef unpack_param(struct si_shader_context *ctx,
 
        if (rshift)
                value = LLVMBuildLShr(gallivm->builder, value,
-                                     lp_build_const_int32(gallivm, rshift), "");
+                                     LLVMConstInt(ctx->i32, rshift, 0), "");
 
        if (rshift + bitwidth < 32) {
                unsigned mask = (1 << bitwidth) - 1;
                value = LLVMBuildAnd(gallivm->builder, value,
-                                    lp_build_const_int32(gallivm, mask), "");
+                                    LLVMConstInt(ctx->i32, mask, 0), "");
        }
 
        return value;
@@ -225,9 +227,9 @@ static LLVMValueRef
 get_tcs_in_patch_stride(struct si_shader_context *ctx)
 {
        if (ctx->type == PIPE_SHADER_VERTEX)
-               return unpack_param(ctx, SI_PARAM_LS_OUT_LAYOUT, 0, 13);
+               return unpack_param(ctx, SI_PARAM_LS_OUT_LAYOUT, 8, 13);
        else if (ctx->type == PIPE_SHADER_TESS_CTRL)
-               return unpack_param(ctx, SI_PARAM_TCS_IN_LAYOUT, 0, 13);
+               return unpack_param(ctx, SI_PARAM_TCS_IN_LAYOUT, 8, 13);
        else {
                assert(0);
                return NULL;
@@ -300,23 +302,21 @@ get_tcs_out_current_patch_data_offset(struct si_shader_context *ctx)
 }
 
 static LLVMValueRef get_instance_index_for_fetch(
-       struct si_shader_context *radeon_bld,
+       struct si_shader_context *ctx,
        unsigned param_start_instance, unsigned divisor)
 {
-       struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->bld_base);
-       struct gallivm_state *gallivm = radeon_bld->bld_base.base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
 
-       LLVMValueRef result = LLVMGetParam(radeon_bld->main_fn,
+       LLVMValueRef result = LLVMGetParam(ctx->main_fn,
                                           ctx->param_instance_id);
 
        /* The division must be done before START_INSTANCE is added. */
        if (divisor > 1)
                result = LLVMBuildUDiv(gallivm->builder, result,
-                               lp_build_const_int32(gallivm, divisor), "");
+                               LLVMConstInt(ctx->i32, divisor, 0), "");
 
        return LLVMBuildAdd(gallivm->builder, result,
-                           LLVMGetParam(radeon_bld->main_fn, param_start_instance), "");
+                           LLVMGetParam(ctx->main_fn, param_start_instance), "");
 }
 
 /* Bitcast <4 x float> to <2 x double>, extract the component, and convert
@@ -340,8 +340,7 @@ static void declare_input_vs(
        const struct tgsi_full_declaration *decl,
        LLVMValueRef out[4])
 {
-       struct lp_build_context *base = &ctx->bld_base.base;
-       struct gallivm_state *gallivm = base->gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
 
        unsigned chan;
        unsigned fix_fetch;
@@ -357,7 +356,7 @@ static void declare_input_vs(
        /* 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_offset = LLVMConstInt(ctx->i32, input_index, 0);
 
        t_list = ac_build_indexed_load_const(&ctx->ac, t_list_ptr, t_offset);
 
@@ -402,7 +401,7 @@ static void declare_input_vs(
 
        /* Break up the vec4 into individual components */
        for (chan = 0; chan < 4; chan++) {
-               LLVMValueRef llvm_chan = lp_build_const_int32(gallivm, chan);
+               LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, 0);
                out[chan] = LLVMBuildExtractElement(gallivm->builder,
                                                    input[0], llvm_chan, "");
        }
@@ -547,7 +546,7 @@ static LLVMValueRef get_primitive_id(struct lp_build_tgsi_context *bld_base,
        struct si_shader_context *ctx = si_shader_context(bld_base);
 
        if (swizzle > 0)
-               return bld_base->uint_bld.zero;
+               return ctx->i32_0;
 
        switch (ctx->type) {
        case PIPE_SHADER_VERTEX:
@@ -564,7 +563,7 @@ static LLVMValueRef get_primitive_id(struct lp_build_tgsi_context *bld_base,
                                    SI_PARAM_PRIMITIVE_ID);
        default:
                assert(0);
-               return bld_base->uint_bld.zero;
+               return ctx->i32_0;
        }
 }
 
@@ -576,13 +575,13 @@ static LLVMValueRef get_indirect_index(struct si_shader_context *ctx,
                                       const struct tgsi_ind_register *ind,
                                       int rel_index)
 {
-       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef result;
 
        result = ctx->addrs[ind->Index][ind->Swizzle];
        result = LLVMBuildLoad(gallivm->builder, result, "");
        result = LLVMBuildAdd(gallivm->builder, result,
-                             lp_build_const_int32(gallivm, rel_index), "");
+                             LLVMConstInt(ctx->i32, rel_index, 0), "");
        return result;
 }
 
@@ -616,7 +615,7 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
                                   LLVMValueRef vertex_dw_stride,
                                   LLVMValueRef base_addr)
 {
-       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        struct tgsi_shader_info *info = &ctx->shader->selector->info;
        ubyte *name, *index, *array_first;
        int first, param;
@@ -644,7 +643,7 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
                        index = get_indirect_index(ctx, &reg.DimIndirect,
                                                   reg.Dimension.Index);
                else
-                       index = lp_build_const_int32(gallivm, reg.Dimension.Index);
+                       index = LLVMConstInt(ctx->i32, reg.Dimension.Index, 0);
 
                base_addr = LLVMBuildAdd(gallivm->builder, base_addr,
                                         LLVMBuildMul(gallivm->builder, index,
@@ -679,7 +678,7 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
 
                base_addr = LLVMBuildAdd(gallivm->builder, base_addr,
                                    LLVMBuildMul(gallivm->builder, ind_index,
-                                                lp_build_const_int32(gallivm, 4), ""), "");
+                                                LLVMConstInt(ctx->i32, 4, 0), ""), "");
 
                param = si_shader_io_get_unique_index(name[first], index[first]);
        } else {
@@ -689,7 +688,7 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
 
        /* Add the base address of the element. */
        return LLVMBuildAdd(gallivm->builder, base_addr,
-                           lp_build_const_int32(gallivm, param * 4), "");
+                           LLVMConstInt(ctx->i32, param * 4, 0), "");
 }
 
 /* The offchip buffer layout for TCS->TES is
@@ -715,7 +714,7 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx,
                                                LLVMValueRef vertex_index,
                                                LLVMValueRef param_index)
 {
-       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef base_addr, vertices_per_patch, num_patches, total_vertices;
        LLVMValueRef param_stride, constant16;
 
@@ -724,7 +723,7 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx,
        total_vertices = LLVMBuildMul(gallivm->builder, vertices_per_patch,
                                      num_patches, "");
 
-       constant16 = lp_build_const_int32(gallivm, 16);
+       constant16 = LLVMConstInt(ctx->i32, 16, 0);
        if (vertex_index) {
                base_addr = LLVMBuildMul(gallivm->builder, rel_patch_id,
                                         vertices_per_patch, "");
@@ -759,7 +758,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->bld_base.base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        struct tgsi_shader_info *info = &ctx->shader->selector->info;
        ubyte *name, *index, *array_first;
        struct tgsi_full_src_register reg;
@@ -775,8 +774,7 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
                        vertex_index = get_indirect_index(ctx, &reg.DimIndirect,
                                                          reg.Dimension.Index);
                else
-                       vertex_index = lp_build_const_int32(gallivm,
-                                                           reg.Dimension.Index);
+                       vertex_index = LLVMConstInt(ctx->i32, reg.Dimension.Index, 0);
        }
 
        /* Get information about the register. */
@@ -804,14 +802,14 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
 
        } else {
                param_base = reg.Register.Index;
-               param_index = lp_build_const_int32(gallivm, 0);
+               param_index = ctx->i32_0;
        }
 
        param_index_base = si_shader_io_get_unique_index(name[param_base],
                                                         index[param_base]);
 
        param_index = LLVMBuildAdd(gallivm->builder, param_index,
-                                  lp_build_const_int32(gallivm, param_index_base),
+                                  LLVMConstInt(ctx->i32, param_index_base, 0),
                                   "");
 
        return get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx),
@@ -824,7 +822,7 @@ static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
                                 LLVMValueRef base, bool readonly_memory)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef value, value2;
        LLVMTypeRef llvm_type = tgsi2llvmtype(bld_base, type);
        LLVMTypeRef vec_type = LLVMVectorType(llvm_type, 4);
@@ -842,7 +840,7 @@ static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
 
                value = LLVMBuildBitCast(gallivm->builder, value, vec_type, "");
                return LLVMBuildExtractElement(gallivm->builder, value,
-                                   lp_build_const_int32(gallivm, swizzle), "");
+                                   LLVMConstInt(ctx->i32, swizzle, 0), "");
        }
 
        value = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
@@ -866,7 +864,7 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
                             LLVMValueRef dw_addr)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef value;
 
        if (swizzle == ~0) {
@@ -875,18 +873,18 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
                for (unsigned chan = 0; chan < TGSI_NUM_CHANNELS; chan++)
                        values[chan] = lds_load(bld_base, type, chan, dw_addr);
 
-               return lp_build_gather_values(bld_base->base.gallivm, values,
+               return lp_build_gather_values(gallivm, values,
                                              TGSI_NUM_CHANNELS);
        }
 
        dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
-                           lp_build_const_int32(gallivm, swizzle));
+                           LLVMConstInt(ctx->i32, swizzle, 0));
 
        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));
+                                      ctx->i32_1);
                value2 = ac_build_indexed_load(&ctx->ac, ctx->lds, dw_addr, false);
                return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
        }
@@ -907,10 +905,10 @@ static void lds_store(struct lp_build_tgsi_context *bld_base,
                      LLVMValueRef value)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
 
        dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
-                           lp_build_const_int32(gallivm, swizzle));
+                           LLVMConstInt(ctx->i32, swizzle, 0));
 
        value = LLVMBuildBitCast(gallivm->builder, value, ctx->i32, "");
        ac_build_indexed_store(&ctx->ac, ctx->lds,
@@ -925,7 +923,7 @@ static LLVMValueRef fetch_input_tcs(
        struct si_shader_context *ctx = si_shader_context(bld_base);
        LLVMValueRef dw_addr, stride;
 
-       stride = unpack_param(ctx, SI_PARAM_TCS_IN_LAYOUT, 13, 8);
+       stride = unpack_param(ctx, SI_PARAM_TCS_IN_LAYOUT, 24, 8);
        dw_addr = get_tcs_in_current_patch_offset(ctx);
        dw_addr = get_dw_address(ctx, NULL, reg, stride, dw_addr);
 
@@ -958,13 +956,12 @@ static LLVMValueRef fetch_input_tes(
        enum tgsi_opcode_type type, unsigned swizzle)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
        LLVMValueRef rw_buffers, buffer, base, addr;
 
        rw_buffers = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_RW_BUFFERS);
        buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
+                       LLVMConstInt(ctx->i32, SI_HS_RING_TESS_OFFCHIP, 0));
 
        base = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
        addr = get_tcs_tes_buffer_address_from_reg(ctx, NULL, reg);
@@ -978,7 +975,7 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
                             LLVMValueRef dst[4])
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        const struct tgsi_full_dst_register *reg = &inst->Dst[0];
        const struct tgsi_shader_info *sh_info = &ctx->shader->selector->info;
        unsigned chan_index;
@@ -1022,7 +1019,7 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
        rw_buffers = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_RW_BUFFERS);
        buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
+                       LLVMConstInt(ctx->i32, SI_HS_RING_TESS_OFFCHIP, 0));
 
        base = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
        buf_addr = get_tcs_tes_buffer_address_from_reg(ctx, reg, NULL);
@@ -1049,7 +1046,7 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
        }
 
        if (inst->Dst[0].Register.WriteMask == 0xF && !is_tess_factor) {
-               LLVMValueRef value = lp_build_gather_values(bld_base->base.gallivm,
+               LLVMValueRef value = lp_build_gather_values(gallivm,
                                                            values, 4);
                ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buf_addr,
                                            base, 0, 1, 0, true, false);
@@ -1062,11 +1059,10 @@ static LLVMValueRef fetch_input_gs(
        enum tgsi_opcode_type type,
        unsigned swizzle)
 {
-       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->bld_base.uint_bld;
-       struct gallivm_state *gallivm = base->gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef vtx_offset, soffset;
        unsigned vtx_offset_param;
        struct tgsi_shader_info *info = &shader->selector->info;
@@ -1087,7 +1083,7 @@ static LLVMValueRef fetch_input_gs(
                for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
                        values[chan] = fetch_input_gs(bld_base, reg, type, chan);
                }
-               return lp_build_gather_values(bld_base->base.gallivm, values,
+               return lp_build_gather_values(gallivm, values,
                                              TGSI_NUM_CHANNELS);
        }
 
@@ -1107,14 +1103,14 @@ static LLVMValueRef fetch_input_gs(
        param = si_shader_io_get_unique_index(semantic_name, semantic_index);
        soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle) * 256, 0);
 
-       value = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1, uint->zero,
+       value = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1, ctx->i32_0,
                                     vtx_offset, soffset, 0, 1, 0, true);
        if (tgsi_type_is_64bit(type)) {
                LLVMValueRef value2;
                soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle + 1) * 256, 0);
 
                value2 = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1,
-                                             uint->zero, vtx_offset, soffset,
+                                             ctx->i32_0, vtx_offset, soffset,
                                              0, 1, 0, true);
                return si_llvm_emit_fetch_64bit(bld_base, type,
                                                value, value2);
@@ -1178,10 +1174,7 @@ static void interp_fs_input(struct si_shader_context *ctx,
                            LLVMValueRef face,
                            LLVMValueRef result[4])
 {
-       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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef attr_number;
        LLVMValueRef i, j;
 
@@ -1202,16 +1195,16 @@ static void interp_fs_input(struct si_shader_context *ctx,
         */
        bool interp = interp_param != NULL;
 
-       attr_number = lp_build_const_int32(gallivm, input_index);
+       attr_number = LLVMConstInt(ctx->i32, input_index, 0);
 
        if (interp) {
                interp_param = LLVMBuildBitCast(gallivm->builder, interp_param,
                                                LLVMVectorType(ctx->f32, 2), "");
 
                i = LLVMBuildExtractElement(gallivm->builder, interp_param,
-                                               uint->zero, "");
+                                               ctx->i32_0, "");
                j = LLVMBuildExtractElement(gallivm->builder, interp_param,
-                                               uint->one, "");
+                                               ctx->i32_1, "");
        }
 
        if (semantic_name == TGSI_SEMANTIC_COLOR &&
@@ -1226,13 +1219,13 @@ static void interp_fs_input(struct si_shader_context *ctx,
                if (semantic_index == 1 && colors_read_mask & 0xf)
                        back_attr_offset += 1;
 
-               back_attr_number = lp_build_const_int32(gallivm, back_attr_offset);
+               back_attr_number = LLVMConstInt(ctx->i32, back_attr_offset, 0);
 
                is_face_positive = LLVMBuildICmp(gallivm->builder, LLVMIntNE,
-                                                face, uint->zero, "");
+                                                face, ctx->i32_0, "");
 
                for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
-                       LLVMValueRef llvm_chan = lp_build_const_int32(gallivm, chan);
+                       LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, 0);
                        LLVMValueRef front, back;
 
                        if (interp) {
@@ -1244,10 +1237,10 @@ static void interp_fs_input(struct si_shader_context *ctx,
                                                        i, j);
                        } else {
                                front = ac_build_fs_interp_mov(&ctx->ac,
-                                       lp_build_const_int32(gallivm, 2), /* P0 */
+                                       LLVMConstInt(ctx->i32, 2, 0), /* P0 */
                                        llvm_chan, attr_number, prim_mask);
                                back = ac_build_fs_interp_mov(&ctx->ac,
-                                       lp_build_const_int32(gallivm, 2), /* P0 */
+                                       LLVMConstInt(ctx->i32, 2, 0), /* P0 */
                                        llvm_chan, back_attr_number, prim_mask);
                        }
 
@@ -1259,26 +1252,26 @@ static void interp_fs_input(struct si_shader_context *ctx,
                }
        } else if (semantic_name == TGSI_SEMANTIC_FOG) {
                if (interp) {
-                       result[0] = ac_build_fs_interp(&ctx->ac, uint->zero,
+                       result[0] = ac_build_fs_interp(&ctx->ac, ctx->i32_0,
                                                       attr_number, prim_mask, i, j);
                } else {
-                       result[0] = ac_build_fs_interp_mov(&ctx->ac, uint->zero,
-                                                          lp_build_const_int32(gallivm, 2), /* P0 */
+                       result[0] = ac_build_fs_interp_mov(&ctx->ac, ctx->i32_0,
+                                                          LLVMConstInt(ctx->i32, 2, 0), /* P0 */
                                                           attr_number, prim_mask);
                }
                result[1] =
-               result[2] = lp_build_const_float(gallivm, 0.0f);
-               result[3] = lp_build_const_float(gallivm, 1.0f);
+               result[2] = LLVMConstReal(ctx->f32, 0.0f);
+               result[3] = LLVMConstReal(ctx->f32, 1.0f);
        } else {
                for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
-                       LLVMValueRef llvm_chan = lp_build_const_int32(gallivm, chan);
+                       LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, 0);
 
                        if (interp) {
                                result[chan] = ac_build_fs_interp(&ctx->ac,
                                        llvm_chan, attr_number, prim_mask, i, j);
                        } else {
                                result[chan] = ac_build_fs_interp_mov(&ctx->ac,
-                                       lp_build_const_int32(gallivm, 2), /* P0 */
+                                       LLVMConstInt(ctx->i32, 2, 0), /* P0 */
                                        llvm_chan, attr_number, prim_mask);
                        }
                }
@@ -1286,16 +1279,14 @@ static void interp_fs_input(struct si_shader_context *ctx,
 }
 
 static void declare_input_fs(
-       struct si_shader_context *radeon_bld,
+       struct si_shader_context *ctx,
        unsigned input_index,
        const struct tgsi_full_declaration *decl,
        LLVMValueRef out[4])
 {
-       struct lp_build_context *base = &radeon_bld->bld_base.base;
-       struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->bld_base);
+       struct lp_build_context *base = &ctx->bld_base.base;
        struct si_shader *shader = ctx->shader;
-       LLVMValueRef main_fn = radeon_bld->main_fn;
+       LLVMValueRef main_fn = ctx->main_fn;
        LLVMValueRef interp_param = NULL;
        int interp_param_idx;
 
@@ -1335,10 +1326,9 @@ static void declare_input_fs(
                        &out[0]);
 }
 
-static LLVMValueRef get_sample_id(struct si_shader_context *radeon_bld)
+static LLVMValueRef get_sample_id(struct si_shader_context *ctx)
 {
-       return unpack_param(si_shader_context(&radeon_bld->bld_base),
-                           SI_PARAM_ANCILLARY, 8, 4);
+       return unpack_param(ctx, SI_PARAM_ANCILLARY, 8, 4);
 }
 
 
@@ -1357,73 +1347,71 @@ static LLVMValueRef buffer_load_const(struct si_shader_context *ctx,
                                  LP_FUNC_ATTR_LEGACY);
 }
 
-static LLVMValueRef load_sample_position(struct si_shader_context *radeon_bld, LLVMValueRef sample_id)
+static LLVMValueRef load_sample_position(struct si_shader_context *ctx, LLVMValueRef sample_id)
 {
-       struct si_shader_context *ctx =
-               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;
+       struct lp_build_context *uint_bld = &ctx->bld_base.uint_bld;
+       struct gallivm_state *gallivm = &ctx->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 buf_index = LLVMConstInt(ctx->i32, SI_PS_CONST_SAMPLE_POSITIONS, 0);
        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);
-       LLVMValueRef offset1 = LLVMBuildAdd(builder, offset0, lp_build_const_int32(gallivm, 4), "");
+       LLVMValueRef offset1 = LLVMBuildAdd(builder, offset0, LLVMConstInt(ctx->i32, 4, 0), "");
 
        LLVMValueRef pos[4] = {
                buffer_load_const(ctx, resource, offset0),
                buffer_load_const(ctx, resource, offset1),
-               lp_build_const_float(gallivm, 0),
-               lp_build_const_float(gallivm, 0)
+               LLVMConstReal(ctx->f32, 0),
+               LLVMConstReal(ctx->f32, 0)
        };
 
        return lp_build_gather_values(gallivm, pos, 4);
 }
 
-static void declare_system_value(
-       struct si_shader_context *radeon_bld,
-       unsigned index,
-       const struct tgsi_full_declaration *decl)
+static void declare_system_value(struct si_shader_context *ctx,
+                                unsigned index,
+                                const struct tgsi_full_declaration *decl)
 {
-       struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->bld_base);
-       struct lp_build_context *bld = &radeon_bld->bld_base.base;
-       struct gallivm_state *gallivm = &radeon_bld->gallivm;
+       struct lp_build_context *bld = &ctx->bld_base.base;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef value = 0;
 
+       assert(index < RADEON_LLVM_MAX_SYSTEM_VALUES);
+
        switch (decl->Semantic.Name) {
        case TGSI_SEMANTIC_INSTANCEID:
-               value = LLVMGetParam(radeon_bld->main_fn,
+               value = LLVMGetParam(ctx->main_fn,
                                     ctx->param_instance_id);
                break;
 
        case TGSI_SEMANTIC_VERTEXID:
                value = LLVMBuildAdd(gallivm->builder,
-                                    LLVMGetParam(radeon_bld->main_fn,
+                                    LLVMGetParam(ctx->main_fn,
                                                  ctx->param_vertex_id),
-                                    LLVMGetParam(radeon_bld->main_fn,
+                                    LLVMGetParam(ctx->main_fn,
                                                  SI_PARAM_BASE_VERTEX), "");
                break;
 
        case TGSI_SEMANTIC_VERTEXID_NOBASE:
-               value = LLVMGetParam(radeon_bld->main_fn,
-                                    ctx->param_vertex_id);
+               /* Unused. Clarify the meaning in indexed vs. non-indexed
+                * draws if this is ever used again. */
+               assert(false);
                break;
 
        case TGSI_SEMANTIC_BASEVERTEX:
-               value = LLVMGetParam(radeon_bld->main_fn,
+               value = LLVMGetParam(ctx->main_fn,
                                     SI_PARAM_BASE_VERTEX);
                break;
 
        case TGSI_SEMANTIC_BASEINSTANCE:
-               value = LLVMGetParam(radeon_bld->main_fn,
+               value = LLVMGetParam(ctx->main_fn,
                                     SI_PARAM_START_INSTANCE);
                break;
 
        case TGSI_SEMANTIC_DRAWID:
-               value = LLVMGetParam(radeon_bld->main_fn,
+               value = LLVMGetParam(ctx->main_fn,
                                     SI_PARAM_DRAWID);
                break;
 
@@ -1431,7 +1419,7 @@ static void declare_system_value(
                if (ctx->type == PIPE_SHADER_TESS_CTRL)
                        value = unpack_param(ctx, SI_PARAM_REL_IDS, 8, 5);
                else if (ctx->type == PIPE_SHADER_GEOMETRY)
-                       value = LLVMGetParam(radeon_bld->main_fn,
+                       value = LLVMGetParam(ctx->main_fn,
                                             SI_PARAM_GS_INSTANCE_ID);
                else
                        assert(!"INVOCATIONID not implemented");
@@ -1440,11 +1428,11 @@ static void declare_system_value(
        case TGSI_SEMANTIC_POSITION:
        {
                LLVMValueRef pos[4] = {
-                       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->bld_base, TGSI_OPCODE_RCP,
-                                                LLVMGetParam(radeon_bld->main_fn,
+                       LLVMGetParam(ctx->main_fn, SI_PARAM_POS_X_FLOAT),
+                       LLVMGetParam(ctx->main_fn, SI_PARAM_POS_Y_FLOAT),
+                       LLVMGetParam(ctx->main_fn, SI_PARAM_POS_Z_FLOAT),
+                       lp_build_emit_llvm_unary(&ctx->bld_base, TGSI_OPCODE_RCP,
+                                                LLVMGetParam(ctx->main_fn,
                                                              SI_PARAM_POS_W_FLOAT)),
                };
                value = lp_build_gather_values(gallivm, pos, 4);
@@ -1452,23 +1440,23 @@ static void declare_system_value(
        }
 
        case TGSI_SEMANTIC_FACE:
-               value = LLVMGetParam(radeon_bld->main_fn, SI_PARAM_FRONT_FACE);
+               value = LLVMGetParam(ctx->main_fn, SI_PARAM_FRONT_FACE);
                break;
 
        case TGSI_SEMANTIC_SAMPLEID:
-               value = get_sample_id(radeon_bld);
+               value = get_sample_id(ctx);
                break;
 
        case TGSI_SEMANTIC_SAMPLEPOS: {
                LLVMValueRef pos[4] = {
-                       LLVMGetParam(radeon_bld->main_fn, SI_PARAM_POS_X_FLOAT),
-                       LLVMGetParam(radeon_bld->main_fn, SI_PARAM_POS_Y_FLOAT),
-                       lp_build_const_float(gallivm, 0),
-                       lp_build_const_float(gallivm, 0)
+                       LLVMGetParam(ctx->main_fn, SI_PARAM_POS_X_FLOAT),
+                       LLVMGetParam(ctx->main_fn, SI_PARAM_POS_Y_FLOAT),
+                       LLVMConstReal(ctx->f32, 0),
+                       LLVMConstReal(ctx->f32, 0)
                };
-               pos[0] = lp_build_emit_llvm_unary(&radeon_bld->bld_base,
+               pos[0] = lp_build_emit_llvm_unary(&ctx->bld_base,
                                                  TGSI_OPCODE_FRC, pos[0]);
-               pos[1] = lp_build_emit_llvm_unary(&radeon_bld->bld_base,
+               pos[1] = lp_build_emit_llvm_unary(&ctx->bld_base,
                                                  TGSI_OPCODE_FRC, pos[1]);
                value = lp_build_gather_values(gallivm, pos, 4);
                break;
@@ -1478,14 +1466,14 @@ static void declare_system_value(
                /* This can only occur with the OpenGL Core profile, which
                 * doesn't support smoothing.
                 */
-               value = LLVMGetParam(radeon_bld->main_fn, SI_PARAM_SAMPLE_COVERAGE);
+               value = LLVMGetParam(ctx->main_fn, SI_PARAM_SAMPLE_COVERAGE);
                break;
 
        case TGSI_SEMANTIC_TESSCOORD:
        {
                LLVMValueRef coord[4] = {
-                       LLVMGetParam(radeon_bld->main_fn, ctx->param_tes_u),
-                       LLVMGetParam(radeon_bld->main_fn, ctx->param_tes_v),
+                       LLVMGetParam(ctx->main_fn, ctx->param_tes_u),
+                       LLVMGetParam(ctx->main_fn, ctx->param_tes_v),
                        bld->zero,
                        bld->zero
                };
@@ -1518,13 +1506,13 @@ static void declare_system_value(
                rw_buffers = LLVMGetParam(ctx->main_fn,
                                        SI_PARAM_RW_BUFFERS);
                buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
+                       LLVMConstInt(ctx->i32, SI_HS_RING_TESS_OFFCHIP, 0));
 
                base = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
                addr = get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx), NULL,
-                                         lp_build_const_int32(gallivm, param));
+                                         LLVMConstInt(ctx->i32, param, 0));
 
-               value = buffer_load(&radeon_bld->bld_base, TGSI_TYPE_FLOAT,
+               value = buffer_load(&ctx->bld_base, TGSI_TYPE_FLOAT,
                                    ~0, buffer, base, addr, true);
 
                break;
@@ -1536,24 +1524,24 @@ static void declare_system_value(
                LLVMValueRef buf, slot, val[4];
                int i, offset;
 
-               slot = lp_build_const_int32(gallivm, SI_HS_CONST_DEFAULT_TESS_LEVELS);
+               slot = LLVMConstInt(ctx->i32, SI_HS_CONST_DEFAULT_TESS_LEVELS, 0);
                buf = LLVMGetParam(ctx->main_fn, SI_PARAM_RW_BUFFERS);
                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++)
                        val[i] = buffer_load_const(ctx, buf,
-                                                  lp_build_const_int32(gallivm, (offset + i) * 4));
+                                                  LLVMConstInt(ctx->i32, (offset + i) * 4, 0));
                value = lp_build_gather_values(gallivm, val, 4);
                break;
        }
 
        case TGSI_SEMANTIC_PRIMID:
-               value = get_primitive_id(&radeon_bld->bld_base, 0);
+               value = get_primitive_id(&ctx->bld_base, 0);
                break;
 
        case TGSI_SEMANTIC_GRID_SIZE:
-               value = LLVMGetParam(radeon_bld->main_fn, SI_PARAM_GRID_SIZE);
+               value = LLVMGetParam(ctx->main_fn, SI_PARAM_GRID_SIZE);
                break;
 
        case TGSI_SEMANTIC_BLOCK_SIZE:
@@ -1570,21 +1558,21 @@ static void declare_system_value(
                        };
 
                        for (i = 0; i < 3; ++i)
-                               values[i] = lp_build_const_int32(gallivm, sizes[i]);
+                               values[i] = LLVMConstInt(ctx->i32, sizes[i], 0);
 
                        value = lp_build_gather_values(gallivm, values, 3);
                } else {
-                       value = LLVMGetParam(radeon_bld->main_fn, SI_PARAM_BLOCK_SIZE);
+                       value = LLVMGetParam(ctx->main_fn, SI_PARAM_BLOCK_SIZE);
                }
                break;
        }
 
        case TGSI_SEMANTIC_BLOCK_ID:
-               value = LLVMGetParam(radeon_bld->main_fn, SI_PARAM_BLOCK_ID);
+               value = LLVMGetParam(ctx->main_fn, SI_PARAM_BLOCK_ID);
                break;
 
        case TGSI_SEMANTIC_THREAD_ID:
-               value = LLVMGetParam(radeon_bld->main_fn, SI_PARAM_THREAD_ID);
+               value = LLVMGetParam(ctx->main_fn, SI_PARAM_THREAD_ID);
                break;
 
        case TGSI_SEMANTIC_HELPER_INVOCATION:
@@ -1601,21 +1589,59 @@ static void declare_system_value(
                }
                break;
 
+       case TGSI_SEMANTIC_SUBGROUP_SIZE:
+               value = LLVMConstInt(ctx->i32, 64, 0);
+               break;
+
+       case TGSI_SEMANTIC_SUBGROUP_INVOCATION:
+               value = ac_get_thread_id(&ctx->ac);
+               break;
+
+       case TGSI_SEMANTIC_SUBGROUP_EQ_MASK:
+       {
+               LLVMValueRef id = ac_get_thread_id(&ctx->ac);
+               id = LLVMBuildZExt(gallivm->builder, id, ctx->i64, "");
+               value = LLVMBuildShl(gallivm->builder, LLVMConstInt(ctx->i64, 1, 0), id, "");
+               value = LLVMBuildBitCast(gallivm->builder, value, ctx->v2i32, "");
+               break;
+       }
+
+       case TGSI_SEMANTIC_SUBGROUP_GE_MASK:
+       case TGSI_SEMANTIC_SUBGROUP_GT_MASK:
+       case TGSI_SEMANTIC_SUBGROUP_LE_MASK:
+       case TGSI_SEMANTIC_SUBGROUP_LT_MASK:
+       {
+               LLVMValueRef id = ac_get_thread_id(&ctx->ac);
+               if (decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_GT_MASK ||
+                   decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_LE_MASK) {
+                       /* All bits set except LSB */
+                       value = LLVMConstInt(ctx->i64, -2, 0);
+               } else {
+                       /* All bits set */
+                       value = LLVMConstInt(ctx->i64, -1, 0);
+               }
+               id = LLVMBuildZExt(gallivm->builder, id, ctx->i64, "");
+               value = LLVMBuildShl(gallivm->builder, value, id, "");
+               if (decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_LE_MASK ||
+                   decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_LT_MASK)
+                       value = LLVMBuildNot(gallivm->builder, value, "");
+               value = LLVMBuildBitCast(gallivm->builder, value, ctx->v2i32, "");
+               break;
+       }
+
        default:
                assert(!"unknown system value");
                return;
        }
 
-       radeon_bld->system_values[index] = value;
+       ctx->system_values[index] = value;
 }
 
-static void declare_compute_memory(struct si_shader_context *radeon_bld,
+static void declare_compute_memory(struct si_shader_context *ctx,
                                    const struct tgsi_full_declaration *decl)
 {
-       struct si_shader_context *ctx =
-               si_shader_context(&radeon_bld->bld_base);
        struct si_shader_selector *sel = ctx->shader->selector;
-       struct gallivm_state *gallivm = &radeon_bld->gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
 
        LLVMTypeRef i8p = LLVMPointerType(ctx->i8, LOCAL_ADDR_SPACE);
        LLVMValueRef var;
@@ -1662,7 +1688,7 @@ static LLVMValueRef fetch_constant(
                for (chan = 0; chan < TGSI_NUM_CHANNELS; ++chan)
                        values[chan] = fetch_constant(bld_base, reg, type, chan);
 
-               return lp_build_gather_values(bld_base->base.gallivm, values, 4);
+               return lp_build_gather_values(&ctx->gallivm, values, 4);
        }
 
        buf = reg->Register.Dimension ? reg->Dimension.Index : 0;
@@ -1683,7 +1709,7 @@ static LLVMValueRef fetch_constant(
                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,
-                                   lp_build_const_int32(base->gallivm, idx * 4));
+                                   LLVMConstInt(ctx->i32, idx * 4, 0));
        } else {
                addr = LLVMConstInt(ctx->i32, idx * 4, 0);
        }
@@ -1706,25 +1732,25 @@ static LLVMValueRef fetch_constant(
 }
 
 /* Upper 16 bits must be zero. */
-static LLVMValueRef si_llvm_pack_two_int16(struct gallivm_state *gallivm,
+static LLVMValueRef si_llvm_pack_two_int16(struct si_shader_context *ctx,
                                           LLVMValueRef val[2])
 {
-       return LLVMBuildOr(gallivm->builder, val[0],
-                          LLVMBuildShl(gallivm->builder, val[1],
-                                       lp_build_const_int32(gallivm, 16),
+       return LLVMBuildOr(ctx->gallivm.builder, val[0],
+                          LLVMBuildShl(ctx->gallivm.builder, val[1],
+                                       LLVMConstInt(ctx->i32, 16, 0),
                                        ""), "");
 }
 
 /* Upper 16 bits are ignored and will be dropped. */
-static LLVMValueRef si_llvm_pack_two_int32_as_int16(struct gallivm_state *gallivm,
+static LLVMValueRef si_llvm_pack_two_int32_as_int16(struct si_shader_context *ctx,
                                                    LLVMValueRef val[2])
 {
        LLVMValueRef v[2] = {
-               LLVMBuildAnd(gallivm->builder, val[0],
-                            lp_build_const_int32(gallivm, 0xffff), ""),
+               LLVMBuildAnd(ctx->gallivm.builder, val[0],
+                            LLVMConstInt(ctx->i32, 0xffff, 0), ""),
                val[1],
        };
-       return si_llvm_pack_two_int16(gallivm, v);
+       return si_llvm_pack_two_int16(ctx, v);
 }
 
 /* Initialize arguments for the shader export intrinsic */
@@ -1735,8 +1761,7 @@ static void si_llvm_init_export_args(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 gallivm_state *gallivm = base->gallivm;
-       LLVMBuilderRef builder = base->gallivm->builder;
+       LLVMBuilderRef builder = ctx->gallivm.builder;
        LLVMValueRef val[4];
        unsigned spi_shader_col_format = V_028714_SPI_SHADER_32_ABGR;
        unsigned chan;
@@ -1806,7 +1831,7 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
 
                        packed = ac_build_cvt_pkrtz_f16(&ctx->ac, pack_args);
                        args->out[chan] =
-                               LLVMBuildBitCast(base->gallivm->builder,
+                               LLVMBuildBitCast(ctx->gallivm.builder,
                                                 packed, ctx->f32, "");
                }
                break;
@@ -1815,18 +1840,18 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
                for (chan = 0; chan < 4; chan++) {
                        val[chan] = ac_build_clamp(&ctx->ac, values[chan]);
                        val[chan] = LLVMBuildFMul(builder, val[chan],
-                                                 lp_build_const_float(gallivm, 65535), "");
+                                                 LLVMConstReal(ctx->f32, 65535), "");
                        val[chan] = LLVMBuildFAdd(builder, val[chan],
-                                                 lp_build_const_float(gallivm, 0.5), "");
+                                                 LLVMConstReal(ctx->f32, 0.5), "");
                        val[chan] = LLVMBuildFPToUI(builder, val[chan],
                                                    ctx->i32, "");
                }
 
                args->compr = 1; /* COMPR flag */
                args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int16(gallivm, val));
+                                 si_llvm_pack_two_int16(ctx, val));
                args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int16(gallivm, val+2));
+                                 si_llvm_pack_two_int16(ctx, val+2));
                break;
 
        case V_028714_SPI_SHADER_SNORM16_ABGR:
@@ -1834,35 +1859,35 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
                        /* Clamp between [-1, 1]. */
                        val[chan] = lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_MIN,
                                                              values[chan],
-                                                             lp_build_const_float(gallivm, 1));
+                                                             LLVMConstReal(ctx->f32, 1));
                        val[chan] = lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_MAX,
                                                              val[chan],
-                                                             lp_build_const_float(gallivm, -1));
+                                                             LLVMConstReal(ctx->f32, -1));
                        /* Convert to a signed integer in [-32767, 32767]. */
                        val[chan] = LLVMBuildFMul(builder, val[chan],
-                                                 lp_build_const_float(gallivm, 32767), "");
+                                                 LLVMConstReal(ctx->f32, 32767), "");
                        /* If positive, add 0.5, else add -0.5. */
                        val[chan] = LLVMBuildFAdd(builder, val[chan],
                                        LLVMBuildSelect(builder,
                                                LLVMBuildFCmp(builder, LLVMRealOGE,
                                                              val[chan], base->zero, ""),
-                                               lp_build_const_float(gallivm, 0.5),
-                                               lp_build_const_float(gallivm, -0.5), ""), "");
+                                               LLVMConstReal(ctx->f32, 0.5),
+                                               LLVMConstReal(ctx->f32, -0.5), ""), "");
                        val[chan] = LLVMBuildFPToSI(builder, val[chan], ctx->i32, "");
                }
 
                args->compr = 1; /* COMPR flag */
                args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int32_as_int16(gallivm, val));
+                                 si_llvm_pack_two_int32_as_int16(ctx, val));
                args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int32_as_int16(gallivm, val+2));
+                                 si_llvm_pack_two_int32_as_int16(ctx, val+2));
                break;
 
        case V_028714_SPI_SHADER_UINT16_ABGR: {
-               LLVMValueRef max_rgb = lp_build_const_int32(gallivm,
-                       is_int8 ? 255 : is_int10 ? 1023 : 65535);
+               LLVMValueRef max_rgb = LLVMConstInt(ctx->i32,
+                       is_int8 ? 255 : is_int10 ? 1023 : 65535, 0);
                LLVMValueRef max_alpha =
-                       !is_int10 ? max_rgb : lp_build_const_int32(gallivm, 3);
+                       !is_int10 ? max_rgb : LLVMConstInt(ctx->i32, 3, 0);
 
                /* Clamp. */
                for (chan = 0; chan < 4; chan++) {
@@ -1874,21 +1899,21 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
 
                args->compr = 1; /* COMPR flag */
                args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int16(gallivm, val));
+                                 si_llvm_pack_two_int16(ctx, val));
                args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int16(gallivm, val+2));
+                                 si_llvm_pack_two_int16(ctx, val+2));
                break;
        }
 
        case V_028714_SPI_SHADER_SINT16_ABGR: {
-               LLVMValueRef max_rgb = lp_build_const_int32(gallivm,
-                       is_int8 ? 127 : is_int10 ? 511 : 32767);
-               LLVMValueRef min_rgb = lp_build_const_int32(gallivm,
-                       is_int8 ? -128 : is_int10 ? -512 : -32768);
+               LLVMValueRef max_rgb = LLVMConstInt(ctx->i32,
+                       is_int8 ? 127 : is_int10 ? 511 : 32767, 0);
+               LLVMValueRef min_rgb = LLVMConstInt(ctx->i32,
+                       is_int8 ? -128 : is_int10 ? -512 : -32768, 0);
                LLVMValueRef max_alpha =
-                       !is_int10 ? max_rgb : lp_build_const_int32(gallivm, 1);
+                       !is_int10 ? max_rgb : ctx->i32_1;
                LLVMValueRef min_alpha =
-                       !is_int10 ? min_rgb : lp_build_const_int32(gallivm, -2);
+                       !is_int10 ? min_rgb : LLVMConstInt(ctx->i32, -2, 0);
 
                /* Clamp. */
                for (chan = 0; chan < 4; chan++) {
@@ -1903,9 +1928,9 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
 
                args->compr = 1; /* COMPR flag */
                args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int32_as_int16(gallivm, val));
+                                 si_llvm_pack_two_int32_as_int16(ctx, val));
                args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int32_as_int16(gallivm, val+2));
+                                 si_llvm_pack_two_int32_as_int16(ctx, val+2));
                break;
        }
 
@@ -1919,7 +1944,6 @@ static void si_alpha_test(struct lp_build_tgsi_context *bld_base,
                          LLVMValueRef alpha)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
 
        if (ctx->shader->key.part.ps.epilog.alpha_func != PIPE_FUNC_NEVER) {
                LLVMValueRef alpha_ref = LLVMGetParam(ctx->main_fn,
@@ -1932,8 +1956,8 @@ static void si_alpha_test(struct lp_build_tgsi_context *bld_base,
                LLVMValueRef arg =
                        lp_build_select(&bld_base->base,
                                        alpha_pass,
-                                       lp_build_const_float(gallivm, 1.0f),
-                                       lp_build_const_float(gallivm, -1.0f));
+                                       LLVMConstReal(ctx->f32, 1.0f),
+                                       LLVMConstReal(ctx->f32, -1.0f));
 
                ac_build_kill(&ctx->ac, arg);
        } else {
@@ -1946,7 +1970,7 @@ static LLVMValueRef si_scale_alpha_by_sample_mask(struct lp_build_tgsi_context *
                                                  unsigned samplemask_param)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef coverage;
 
        /* alpha = alpha * popcount(coverage) / SI_NUM_SMOOTH_AA_SAMPLES */
@@ -1962,7 +1986,7 @@ static LLVMValueRef si_scale_alpha_by_sample_mask(struct lp_build_tgsi_context *
                                   ctx->f32, "");
 
        coverage = LLVMBuildFMul(gallivm->builder, coverage,
-                                lp_build_const_float(gallivm,
+                                LLVMConstReal(ctx->f32,
                                        1.0 / SI_NUM_SMOOTH_AA_SAMPLES), "");
 
        return LLVMBuildFMul(gallivm->builder, alpha, coverage, "");
@@ -1978,8 +2002,8 @@ static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base,
        unsigned const_chan;
        LLVMValueRef base_elt;
        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 constbuf_index = LLVMConstInt(ctx->i32,
+                                                  SI_VS_CONST_CLIP_PLANES, 0);
        LLVMValueRef const_resource = ac_build_indexed_load_const(&ctx->ac, ptr, constbuf_index);
 
        for (reg_index = 0; reg_index < 2; reg_index ++) {
@@ -1988,7 +2012,7 @@ static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base,
                args->out[0] =
                args->out[1] =
                args->out[2] =
-               args->out[3] = lp_build_const_float(base->gallivm, 0.0f);
+               args->out[3] = LLVMConstReal(ctx->f32, 0.0f);
 
                /* Compute dot products of position and user clip plane vectors */
                for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
@@ -2081,7 +2105,7 @@ static void emit_streamout_output(struct si_shader_context *ctx,
        ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf_idx],
                                    vdata, num_comps,
                                    so_write_offsets[buf_idx],
-                                   LLVMConstInt(ctx->i32, 0, 0),
+                                   ctx->i32_0,
                                    stream_out->dst_offset * 4, 1, 1, true, false);
 }
 
@@ -2139,8 +2163,8 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
                        if (!so->stride[i])
                                continue;
 
-                       LLVMValueRef offset = lp_build_const_int32(gallivm,
-                                                                  SI_VS_STREAMOUT_BUF0 + i);
+                       LLVMValueRef offset = LLVMConstInt(ctx->i32,
+                                                          SI_VS_STREAMOUT_BUF0 + i, 0);
 
                        so_buffers[i] = ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
 
@@ -2324,15 +2348,15 @@ handle_semantic:
                if (shader->selector->info.writes_edgeflag) {
                        /* The output is a float, but the hw expects an integer
                         * with the first bit containing the edge flag. */
-                       edgeflag_value = LLVMBuildFPToUI(base->gallivm->builder,
+                       edgeflag_value = LLVMBuildFPToUI(ctx->gallivm.builder,
                                                         edgeflag_value,
                                                         ctx->i32, "");
                        edgeflag_value = lp_build_min(&bld_base->int_bld,
                                                      edgeflag_value,
-                                                     bld_base->int_bld.one);
+                                                     ctx->i32_1);
 
                        /* The LLVM intrinsic expects a float. */
-                       pos_args[1].out[1] = LLVMBuildBitCast(base->gallivm->builder,
+                       pos_args[1].out[1] = LLVMBuildBitCast(ctx->gallivm.builder,
                                                          edgeflag_value,
                                                          ctx->f32, "");
                }
@@ -2371,7 +2395,7 @@ handle_semantic:
 static void si_copy_tcs_inputs(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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef invocation_id, rw_buffers, buffer, buffer_offset;
        LLVMValueRef lds_vertex_stride, lds_vertex_offset, lds_base;
        uint64_t inputs;
@@ -2380,11 +2404,11 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
 
        rw_buffers = LLVMGetParam(ctx->main_fn, SI_PARAM_RW_BUFFERS);
        buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
+                       LLVMConstInt(ctx->i32, SI_HS_RING_TESS_OFFCHIP, 0));
 
        buffer_offset = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
 
-       lds_vertex_stride = unpack_param(ctx, SI_PARAM_TCS_IN_LAYOUT, 13, 8);
+       lds_vertex_stride = unpack_param(ctx, SI_PARAM_TCS_IN_LAYOUT, 24, 8);
        lds_vertex_offset = LLVMBuildMul(gallivm->builder, invocation_id,
                                         lds_vertex_stride, "");
        lds_base = get_tcs_in_current_patch_offset(ctx);
@@ -2395,13 +2419,13 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
                unsigned i = u_bit_scan64(&inputs);
 
                LLVMValueRef lds_ptr = LLVMBuildAdd(gallivm->builder, lds_base,
-                                           lp_build_const_int32(gallivm, 4 * i),
+                                           LLVMConstInt(ctx->i32, 4 * i, 0),
                                             "");
 
                LLVMValueRef buffer_addr = get_tcs_tes_buffer_address(ctx,
                                              get_rel_patch_id(ctx),
                                              invocation_id,
-                                             lp_build_const_int32(gallivm, i));
+                                             LLVMConstInt(ctx->i32, i, 0));
 
                LLVMValueRef value = lds_load(bld_base, TGSI_TYPE_SIGNED, ~0,
                                              lds_ptr);
@@ -2417,7 +2441,7 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
                                  LLVMValueRef tcs_out_current_patch_data_offset)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        struct si_shader *shader = ctx->shader;
        unsigned tess_inner_index, tess_outer_index;
        LLVMValueRef lds_base, lds_inner, lds_outer, byteoffset, buffer;
@@ -2435,7 +2459,7 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
         */
        lp_build_if(&if_ctx, gallivm,
                    LLVMBuildICmp(gallivm->builder, LLVMIntEQ,
-                                 invocation_id, bld_base->uint_bld.zero, ""));
+                                 invocation_id, ctx->i32_0, ""));
 
        /* Determine the layout of one tess factor element in the buffer. */
        switch (shader->key.part.tcs.epilog.prim_mode) {
@@ -2467,11 +2491,11 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
 
        lds_base = tcs_out_current_patch_data_offset;
        lds_inner = LLVMBuildAdd(gallivm->builder, lds_base,
-                                lp_build_const_int32(gallivm,
-                                                     tess_inner_index * 4), "");
+                                LLVMConstInt(ctx->i32,
+                                             tess_inner_index * 4, 0), "");
        lds_outer = LLVMBuildAdd(gallivm->builder, lds_base,
-                                lp_build_const_int32(gallivm,
-                                                     tess_outer_index * 4), "");
+                                LLVMConstInt(ctx->i32,
+                                             tess_outer_index * 4, 0), "");
 
        for (i = 0; i < 4; i++) {
                inner[i] = LLVMGetUndef(ctx->i32);
@@ -2506,22 +2530,22 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
        rw_buffers = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_RW_BUFFERS);
        buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_FACTOR));
+                       LLVMConstInt(ctx->i32, SI_HS_RING_TESS_FACTOR, 0));
 
        /* Get the offset. */
        tf_base = LLVMGetParam(ctx->main_fn,
                               SI_PARAM_TESS_FACTOR_OFFSET);
        byteoffset = LLVMBuildMul(gallivm->builder, rel_patch_id,
-                                 lp_build_const_int32(gallivm, 4 * stride), "");
+                                 LLVMConstInt(ctx->i32, 4 * stride, 0), "");
 
        lp_build_if(&inner_if_ctx, gallivm,
                    LLVMBuildICmp(gallivm->builder, LLVMIntEQ,
-                                 rel_patch_id, bld_base->uint_bld.zero, ""));
+                                 rel_patch_id, ctx->i32_0, ""));
 
        /* Store the dynamic HS control word. */
        ac_build_buffer_store_dword(&ctx->ac, buffer,
-                                   lp_build_const_int32(gallivm, 0x80000000),
-                                   1, lp_build_const_int32(gallivm, 0), tf_base,
+                                   LLVMConstInt(ctx->i32, 0x80000000, 0),
+                                   1, ctx->i32_0, tf_base,
                                    0, 1, 0, true, false);
 
        lp_build_endif(&inner_if_ctx);
@@ -2587,7 +2611,7 @@ static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base)
        tf_lds_offset = get_tcs_out_current_patch_data_offset(ctx);
 
        /* Return epilog parameters from this function. */
-       LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+       LLVMBuilderRef builder = ctx->gallivm.builder;
        LLVMValueRef ret = ctx->return_value;
        LLVMValueRef rw_buffers, rw0, rw1, tf_soffset;
        unsigned vgpr;
@@ -2598,9 +2622,9 @@ static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base)
        rw_buffers = LLVMBuildPtrToInt(builder, rw_buffers, ctx->i64, "");
        rw_buffers = LLVMBuildBitCast(builder, rw_buffers, ctx->v2i32, "");
        rw0 = LLVMBuildExtractElement(builder, rw_buffers,
-                                     bld_base->uint_bld.zero, "");
+                                     ctx->i32_0, "");
        rw1 = LLVMBuildExtractElement(builder, rw_buffers,
-                                     bld_base->uint_bld.one, "");
+                                     ctx->i32_1, "");
        ret = LLVMBuildInsertValue(builder, ret, rw0, 0, "");
        ret = LLVMBuildInsertValue(builder, ret, rw1, 1, "");
 
@@ -2634,12 +2658,12 @@ static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base)
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct si_shader *shader = ctx->shader;
        struct tgsi_shader_info *info = &shader->selector->info;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        unsigned i, chan;
        LLVMValueRef vertex_id = LLVMGetParam(ctx->main_fn,
                                              ctx->param_rel_auto_id);
        LLVMValueRef vertex_dw_stride =
-               unpack_param(ctx, SI_PARAM_LS_OUT_LAYOUT, 13, 8);
+               unpack_param(ctx, SI_PARAM_LS_OUT_LAYOUT, 24, 8);
        LLVMValueRef base_dw_addr = LLVMBuildMul(gallivm->builder, vertex_id,
                                                 vertex_dw_stride, "");
 
@@ -2651,7 +2675,7 @@ static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base)
                unsigned index = info->output_semantic_index[i];
                int param = si_shader_io_get_unique_index(name, index);
                LLVMValueRef dw_addr = LLVMBuildAdd(gallivm->builder, base_dw_addr,
-                                       lp_build_const_int32(gallivm, param * 4), "");
+                                       LLVMConstInt(ctx->i32, param * 4, 0), "");
 
                for (chan = 0; chan < 4; chan++) {
                        lds_store(bld_base, chan, dw_addr,
@@ -2663,7 +2687,7 @@ static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base)
 static void si_llvm_emit_es_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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        struct si_shader *es = ctx->shader;
        struct tgsi_shader_info *info = &es->selector->info;
        LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
@@ -2706,7 +2730,7 @@ static void si_llvm_emit_gs_epilogue(struct lp_build_tgsi_context *bld_base)
 static void si_llvm_emit_vs_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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        struct tgsi_shader_info *info = &ctx->shader->selector->info;
        struct si_shader_output_values *outputs = NULL;
        int i,j;
@@ -2839,7 +2863,7 @@ static void si_export_mrt_z(struct lp_build_tgsi_context *bld_base,
                if (stencil) {
                        /* Stencil should be in X[23:16]. */
                        stencil = bitcast(bld_base, TGSI_TYPE_UNSIGNED, stencil);
-                       stencil = LLVMBuildShl(base->gallivm->builder, stencil,
+                       stencil = LLVMBuildShl(ctx->gallivm.builder, stencil,
                                               LLVMConstInt(ctx->i32, 16, 0), "");
                        args.out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT, stencil);
                        mask |= 0x3;
@@ -2987,9 +3011,8 @@ static void si_llvm_return_fs_outputs(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 tgsi_shader_info *info = &shader->selector->info;
-       LLVMBuilderRef builder = base->gallivm->builder;
+       LLVMBuilderRef builder = ctx->gallivm.builder;
        unsigned i, j, first_vgpr, vgpr;
 
        LLVMValueRef color[8][4] = {};
@@ -3073,24 +3096,24 @@ static LLVMValueRef get_buffer_size(
        LLVMValueRef descriptor)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        LLVMValueRef size =
                LLVMBuildExtractElement(builder, descriptor,
-                                       lp_build_const_int32(gallivm, 2), "");
+                                       LLVMConstInt(ctx->i32, 2, 0), "");
 
-       if (ctx->screen->b.chip_class >= VI) {
+       if (ctx->screen->b.chip_class == VI) {
                /* On VI, the descriptor contains the size in bytes,
                 * but TXQ must return the size in elements.
                 * The stride is always non-zero for resources using TXQ.
                 */
                LLVMValueRef stride =
                        LLVMBuildExtractElement(builder, descriptor,
-                                               lp_build_const_int32(gallivm, 1), "");
+                                               ctx->i32_1, "");
                stride = LLVMBuildLShr(builder, stride,
-                                      lp_build_const_int32(gallivm, 16), "");
+                                      LLVMConstInt(ctx->i32, 16, 0), "");
                stride = LLVMBuildAnd(builder, stride,
-                                     lp_build_const_int32(gallivm, 0x3FFF), "");
+                                     LLVMConstInt(ctx->i32, 0x3FFF, 0), "");
 
                size = LLVMBuildUDiv(builder, size, stride, "");
        }
@@ -3105,16 +3128,43 @@ static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
 /* Prevent optimizations (at least of memory accesses) across the current
  * point in the program by emitting empty inline assembly that is marked as
  * having side effects.
+ *
+ * Optionally, a value can be passed through the inline assembly to prevent
+ * LLVM from hoisting calls to ReadNone functions.
  */
-#if 0 /* unused currently */
-static void emit_optimization_barrier(struct si_shader_context *ctx)
+static void emit_optimization_barrier(struct si_shader_context *ctx,
+                                     LLVMValueRef *pvgpr)
 {
+       static int counter = 0;
+
        LLVMBuilderRef builder = ctx->gallivm.builder;
-       LLVMTypeRef ftype = LLVMFunctionType(ctx->voidt, NULL, 0, false);
-       LLVMValueRef inlineasm = LLVMConstInlineAsm(ftype, "", "", true, false);
-       LLVMBuildCall(builder, inlineasm, NULL, 0, "");
+       char code[16];
+
+       snprintf(code, sizeof(code), "; %d", p_atomic_inc_return(&counter));
+
+       if (!pvgpr) {
+               LLVMTypeRef ftype = LLVMFunctionType(ctx->voidt, NULL, 0, false);
+               LLVMValueRef inlineasm = LLVMConstInlineAsm(ftype, code, "", true, false);
+               LLVMBuildCall(builder, inlineasm, NULL, 0, "");
+       } else {
+               LLVMTypeRef ftype = LLVMFunctionType(ctx->i32, &ctx->i32, 1, false);
+               LLVMValueRef inlineasm = LLVMConstInlineAsm(ftype, code, "=v,0", true, false);
+               LLVMValueRef vgpr = *pvgpr;
+               LLVMTypeRef vgpr_type = LLVMTypeOf(vgpr);
+               unsigned vgpr_size = llvm_get_type_size(vgpr_type);
+               LLVMValueRef vgpr0;
+
+               assert(vgpr_size % 4 == 0);
+
+               vgpr = LLVMBuildBitCast(builder, vgpr, LLVMVectorType(ctx->i32, vgpr_size / 4), "");
+               vgpr0 = LLVMBuildExtractElement(builder, vgpr, ctx->i32_0, "");
+               vgpr0 = LLVMBuildCall(builder, inlineasm, &vgpr0, 1, "");
+               vgpr = LLVMBuildInsertElement(builder, vgpr, vgpr0, ctx->i32_0, "");
+               vgpr = LLVMBuildBitCast(builder, vgpr, vgpr_type, "");
+
+               *pvgpr = vgpr;
+       }
 }
-#endif
 
 /* Combine these with & instead of |. */
 #define NOOP_WAITCNT 0xf7f
@@ -3126,7 +3176,7 @@ static void emit_waitcnt(struct si_shader_context *ctx, unsigned simm16)
        struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        LLVMValueRef args[1] = {
-               lp_build_const_int32(gallivm, simm16)
+               LLVMConstInt(ctx->i32, simm16, 0)
        };
        lp_build_intrinsic(builder, "llvm.amdgcn.s.waitcnt",
                           ctx->voidt, args, 1, 0);
@@ -3157,6 +3207,25 @@ static void membar_emit(
                emit_waitcnt(ctx, waitcnt);
 }
 
+static void clock_emit(
+               const struct lp_build_tgsi_action *action,
+               struct lp_build_tgsi_context *bld_base,
+               struct lp_build_emit_data *emit_data)
+{
+       struct si_shader_context *ctx = si_shader_context(bld_base);
+       struct gallivm_state *gallivm = &ctx->gallivm;
+       LLVMValueRef tmp;
+
+       tmp = lp_build_intrinsic(gallivm->builder, "llvm.readcyclecounter",
+                                ctx->i64, NULL, 0, 0);
+       tmp = LLVMBuildBitCast(gallivm->builder, tmp, ctx->v2i32, "");
+
+       emit_data->output[0] =
+               LLVMBuildExtractElement(gallivm->builder, tmp, ctx->i32_0, "");
+       emit_data->output[1] =
+               LLVMBuildExtractElement(gallivm->builder, tmp, ctx->i32_1, "");
+}
+
 static LLVMValueRef
 shader_buffer_fetch_rsrc(struct si_shader_context *ctx,
                         const struct tgsi_full_src_register *reg)
@@ -3230,6 +3299,24 @@ static LLVMTypeRef const_array(LLVMTypeRef elem_type, int num_elements)
                               CONST_ADDR_SPACE);
 }
 
+static LLVMValueRef load_image_desc(struct si_shader_context *ctx,
+                                   LLVMValueRef list, LLVMValueRef index,
+                                   unsigned target)
+{
+       LLVMBuilderRef builder = ctx->gallivm.builder;
+
+       if (target == TGSI_TEXTURE_BUFFER) {
+               index = LLVMBuildMul(builder, index,
+                                    LLVMConstInt(ctx->i32, 2, 0), "");
+               index = LLVMBuildAdd(builder, index,
+                                    ctx->i32_1, "");
+               list = LLVMBuildPointerCast(builder, list,
+                                           const_array(ctx->v4i32, 0), "");
+       }
+
+       return ac_build_indexed_load_const(&ctx->ac, list, index);
+}
+
 /**
  * Load the resource descriptor for \p image.
  */
@@ -3243,8 +3330,8 @@ image_fetch_rsrc(
        struct si_shader_context *ctx = si_shader_context(bld_base);
        LLVMValueRef rsrc_ptr = LLVMGetParam(ctx->main_fn,
                                             SI_PARAM_IMAGES);
-       LLVMValueRef index, tmp;
-       bool dcc_off = target != TGSI_TEXTURE_BUFFER && is_store;
+       LLVMValueRef index;
+       bool dcc_off = is_store;
 
        assert(image->Register.File == TGSI_FILE_IMAGE);
 
@@ -3255,8 +3342,7 @@ image_fetch_rsrc(
 
                index = LLVMConstInt(ctx->i32, image->Register.Index, 0);
 
-               if (images_writemask & (1 << image->Register.Index) &&
-                   target != TGSI_TEXTURE_BUFFER)
+               if (images_writemask & (1 << image->Register.Index))
                        dcc_off = true;
        } else {
                /* From the GL_ARB_shader_image_load_store extension spec:
@@ -3273,23 +3359,9 @@ image_fetch_rsrc(
                                                   SI_NUM_IMAGES);
        }
 
-       if (target == TGSI_TEXTURE_BUFFER) {
-               LLVMBuilderRef builder = ctx->gallivm.builder;
-
-               rsrc_ptr = LLVMBuildPointerCast(builder, rsrc_ptr,
-                                               const_array(ctx->v4i32, 0), "");
-               index = LLVMBuildMul(builder, index,
-                                    LLVMConstInt(ctx->i32, 2, 0), "");
-               index = LLVMBuildAdd(builder, index,
-                                    LLVMConstInt(ctx->i32, 1, 0), "");
-               *rsrc = ac_build_indexed_load_const(&ctx->ac, rsrc_ptr, index);
-               return;
-       }
-
-       tmp = ac_build_indexed_load_const(&ctx->ac, rsrc_ptr, index);
-       if (dcc_off)
-               tmp = force_dcc_off(ctx, tmp);
-       *rsrc = tmp;
+       *rsrc = load_image_desc(ctx, rsrc_ptr, index, target);
+       if (dcc_off && target != TGSI_TEXTURE_BUFFER)
+               *rsrc = force_dcc_off(ctx, *rsrc);
 }
 
 static LLVMValueRef image_fetch_coords(
@@ -3297,7 +3369,8 @@ static LLVMValueRef image_fetch_coords(
                const struct tgsi_full_instruction *inst,
                unsigned src)
 {
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct si_shader_context *ctx = si_shader_context(bld_base);
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        unsigned target = inst->Memory.Texture;
        unsigned num_coords = tgsi_util_get_texture_coord_dim(target);
@@ -3307,10 +3380,21 @@ static LLVMValueRef image_fetch_coords(
 
        for (chan = 0; chan < num_coords; ++chan) {
                tmp = lp_build_emit_fetch(bld_base, inst, src, chan);
-               tmp = LLVMBuildBitCast(builder, tmp, bld_base->uint_bld.elem_type, "");
+               tmp = LLVMBuildBitCast(builder, tmp, ctx->i32, "");
                coords[chan] = tmp;
        }
 
+       /* 1D textures are allocated and used as 2D on GFX9. */
+       if (ctx->screen->b.chip_class >= GFX9) {
+               if (target == TGSI_TEXTURE_1D) {
+                       coords[1] = ctx->i32_0;
+                       num_coords++;
+               } else if (target == TGSI_TEXTURE_1D_ARRAY) {
+                       coords[2] = coords[1];
+                       coords[1] = ctx->i32_0;
+               }
+       }
+
        if (num_coords == 1)
                return coords[0];
 
@@ -3399,12 +3483,12 @@ static void load_fetch_args(
                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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        const struct tgsi_full_instruction * inst = emit_data->inst;
        unsigned target = inst->Memory.Texture;
        LLVMValueRef rsrc;
 
-       emit_data->dst_type = LLVMVectorType(bld_base->base.elem_type, 4);
+       emit_data->dst_type = ctx->v4f32;
 
        if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
                LLVMBuilderRef builder = gallivm->builder;
@@ -3414,9 +3498,9 @@ static void load_fetch_args(
                rsrc = shader_buffer_fetch_rsrc(ctx, &inst->Src[0]);
 
                tmp = lp_build_emit_fetch(bld_base, inst, 1, 0);
-               offset = LLVMBuildBitCast(builder, tmp, bld_base->uint_bld.elem_type, "");
+               offset = LLVMBuildBitCast(builder, tmp, ctx->i32, "");
 
-               buffer_append_args(ctx, emit_data, rsrc, bld_base->uint_bld.zero,
+               buffer_append_args(ctx, emit_data, rsrc, ctx->i32_0,
                                   offset, false, false);
        } else if (inst->Src[0].Register.File == TGSI_FILE_IMAGE) {
                LLVMValueRef coords;
@@ -3426,11 +3510,11 @@ static void load_fetch_args(
 
                if (target == TGSI_TEXTURE_BUFFER) {
                        buffer_append_args(ctx, emit_data, rsrc, coords,
-                                          bld_base->uint_bld.zero, false, false);
+                                          ctx->i32_0, false, false);
                } else {
                        emit_data->args[0] = coords;
                        emit_data->args[1] = rsrc;
-                       emit_data->args[2] = lp_build_const_int32(gallivm, 15); /* dmask */
+                       emit_data->args[2] = LLVMConstInt(ctx->i32, 15, 0); /* dmask */
                        emit_data->arg_count = 3;
 
                        image_append_args(ctx, emit_data, target, false, false);
@@ -3512,54 +3596,27 @@ 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->bld_base.base;
        struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        unsigned writemask = inst->Dst[0].Register.WriteMask;
        LLVMValueRef channels[4], ptr, derived_ptr, index;
        int chan;
 
-       ptr = get_memory_ptr(ctx, inst, base->elem_type, 1);
+       ptr = get_memory_ptr(ctx, inst, ctx->f32, 1);
 
        for (chan = 0; chan < 4; ++chan) {
                if (!(writemask & (1 << chan))) {
-                       channels[chan] = LLVMGetUndef(base->elem_type);
+                       channels[chan] = LLVMGetUndef(ctx->f32);
                        continue;
                }
 
-               index = lp_build_const_int32(gallivm, chan);
+               index = LLVMConstInt(ctx->i32, chan, 0);
                derived_ptr = LLVMBuildGEP(builder, ptr, &index, 1, "");
                channels[chan] = LLVMBuildLoad(builder, derived_ptr, "");
        }
        emit_data->output[emit_data->chan] = lp_build_gather_values(gallivm, channels, 4);
 }
 
-static void get_image_intr_name(const char *base_name,
-                               LLVMTypeRef data_type,
-                               LLVMTypeRef coords_type,
-                               LLVMTypeRef rsrc_type,
-                               char *out_name, unsigned out_len)
-{
-       char coords_type_name[8];
-
-       ac_build_type_name_for_intr(coords_type, coords_type_name,
-                           sizeof(coords_type_name));
-
-       if (HAVE_LLVM <= 0x0309) {
-               snprintf(out_name, out_len, "%s.%s", base_name, coords_type_name);
-       } else {
-               char data_type_name[8];
-               char rsrc_type_name[8];
-
-               ac_build_type_name_for_intr(data_type, data_type_name,
-                                       sizeof(data_type_name));
-               ac_build_type_name_for_intr(rsrc_type, rsrc_type_name,
-                                       sizeof(rsrc_type_name));
-               snprintf(out_name, out_len, "%s.%s.%s.%s", base_name,
-                        data_type_name, coords_type_name, rsrc_type_name);
-       }
-}
-
 /**
  * Return true if the memory accessed by a LOAD or STORE instruction is
  * read-only or write-only, respectively.
@@ -3628,7 +3685,7 @@ static void load_emit(
                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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        const struct tgsi_full_instruction * inst = emit_data->inst;
        const struct tgsi_shader_info *info = &ctx->shader->selector->info;
@@ -3662,11 +3719,11 @@ static void load_emit(
                                emit_data->args, emit_data->arg_count,
                                get_load_intr_attribs(readonly_memory));
        } else {
-               get_image_intr_name("llvm.amdgcn.image.load",
-                               emit_data->dst_type,            /* vdata */
-                               LLVMTypeOf(emit_data->args[0]), /* coords */
-                               LLVMTypeOf(emit_data->args[1]), /* rsrc */
-                               intrinsic_name, sizeof(intrinsic_name));
+               ac_get_image_intr_name("llvm.amdgcn.image.load",
+                                      emit_data->dst_type,             /* vdata */
+                                      LLVMTypeOf(emit_data->args[0]), /* coords */
+                                      LLVMTypeOf(emit_data->args[1]), /* rsrc */
+                                      intrinsic_name, sizeof(intrinsic_name));
 
                emit_data->output[emit_data->chan] =
                        lp_build_intrinsic(
@@ -3681,7 +3738,7 @@ static void store_fetch_args(
                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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        const struct tgsi_full_instruction * inst = emit_data->inst;
        struct tgsi_full_src_register memory;
@@ -3708,9 +3765,9 @@ static void store_fetch_args(
                rsrc = shader_buffer_fetch_rsrc(ctx, &memory);
 
                tmp = lp_build_emit_fetch(bld_base, inst, 0, 0);
-               offset = LLVMBuildBitCast(builder, tmp, bld_base->uint_bld.elem_type, "");
+               offset = LLVMBuildBitCast(builder, tmp, ctx->i32, "");
 
-               buffer_append_args(ctx, emit_data, rsrc, bld_base->uint_bld.zero,
+               buffer_append_args(ctx, emit_data, rsrc, ctx->i32_0,
                                   offset, false, false);
        } else if (inst->Dst[0].Register.File == TGSI_FILE_IMAGE) {
                unsigned target = inst->Memory.Texture;
@@ -3729,12 +3786,12 @@ static void store_fetch_args(
                if (target == TGSI_TEXTURE_BUFFER) {
                        image_fetch_rsrc(bld_base, &memory, true, target, &rsrc);
                        buffer_append_args(ctx, emit_data, rsrc, coords,
-                                          bld_base->uint_bld.zero, false, force_glc);
+                                          ctx->i32_0, false, force_glc);
                } else {
                        emit_data->args[1] = coords;
                        image_fetch_rsrc(bld_base, &memory, true, target,
                                         &emit_data->args[2]);
-                       emit_data->args[3] = lp_build_const_int32(gallivm, 15); /* dmask */
+                       emit_data->args[3] = LLVMConstInt(ctx->i32, 15, 0); /* dmask */
                        emit_data->arg_count = 4;
 
                        image_append_args(ctx, emit_data, target, false, force_glc);
@@ -3750,7 +3807,6 @@ 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->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;
@@ -3779,23 +3835,23 @@ static void store_emit_buffer(
 
                        tmp = LLVMBuildExtractElement(
                                builder, base_data,
-                               lp_build_const_int32(gallivm, start), "");
+                               LLVMConstInt(ctx->i32, start, 0), "");
                        data = LLVMBuildInsertElement(
                                builder, LLVMGetUndef(v2f32), tmp,
-                               uint_bld->zero, "");
+                               ctx->i32_0, "");
 
                        tmp = LLVMBuildExtractElement(
                                builder, base_data,
-                               lp_build_const_int32(gallivm, start + 1), "");
+                               LLVMConstInt(ctx->i32, start + 1, 0), "");
                        data = LLVMBuildInsertElement(
-                               builder, data, tmp, uint_bld->one, "");
+                               builder, data, tmp, ctx->i32_1, "");
 
                        intrinsic_name = "llvm.amdgcn.buffer.store.v2f32";
                } else {
                        assert(count == 1);
                        data = LLVMBuildExtractElement(
                                builder, base_data,
-                               lp_build_const_int32(gallivm, start), "");
+                               LLVMConstInt(ctx->i32, start, 0), "");
                        intrinsic_name = "llvm.amdgcn.buffer.store.f32";
                }
 
@@ -3803,7 +3859,7 @@ static void store_emit_buffer(
                if (start != 0) {
                        offset = LLVMBuildAdd(
                                builder, offset,
-                               lp_build_const_int32(gallivm, start * 4), "");
+                               LLVMConstInt(ctx->i32, start * 4, 0), "");
                }
 
                emit_data->args[0] = data;
@@ -3822,20 +3878,19 @@ 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->bld_base.base;
        LLVMBuilderRef builder = gallivm->builder;
        unsigned writemask = inst->Dst[0].Register.WriteMask;
        LLVMValueRef ptr, derived_ptr, data, index;
        int chan;
 
-       ptr = get_memory_ptr(ctx, inst, base->elem_type, 0);
+       ptr = get_memory_ptr(ctx, inst, ctx->f32, 0);
 
        for (chan = 0; chan < 4; ++chan) {
                if (!(writemask & (1 << chan))) {
                        continue;
                }
                data = lp_build_emit_fetch(&ctx->bld_base, inst, 1, chan);
-               index = lp_build_const_int32(gallivm, chan);
+               index = LLVMConstInt(ctx->i32, chan, 0);
                derived_ptr = LLVMBuildGEP(builder, ptr, &index, 1, "");
                LLVMBuildStore(builder, data, derived_ptr);
        }
@@ -3847,7 +3902,7 @@ static void store_emit(
                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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        const struct tgsi_full_instruction * inst = emit_data->inst;
        const struct tgsi_shader_info *info = &ctx->shader->selector->info;
@@ -3881,11 +3936,11 @@ static void store_emit(
                        emit_data->arg_count,
                        get_store_intr_attribs(writeonly_memory));
        } else {
-               get_image_intr_name("llvm.amdgcn.image.store",
-                               LLVMTypeOf(emit_data->args[0]), /* vdata */
-                               LLVMTypeOf(emit_data->args[1]), /* coords */
-                               LLVMTypeOf(emit_data->args[2]), /* rsrc */
-                               intrinsic_name, sizeof(intrinsic_name));
+               ac_get_image_intr_name("llvm.amdgcn.image.store",
+                                      LLVMTypeOf(emit_data->args[0]), /* vdata */
+                                      LLVMTypeOf(emit_data->args[1]), /* coords */
+                                      LLVMTypeOf(emit_data->args[2]), /* rsrc */
+                                      intrinsic_name, sizeof(intrinsic_name));
 
                emit_data->output[emit_data->chan] =
                        lp_build_intrinsic(
@@ -3900,21 +3955,21 @@ static void atomic_fetch_args(
                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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        const struct tgsi_full_instruction * inst = emit_data->inst;
        LLVMValueRef data1, data2;
        LLVMValueRef rsrc;
        LLVMValueRef tmp;
 
-       emit_data->dst_type = bld_base->base.elem_type;
+       emit_data->dst_type = ctx->f32;
 
        tmp = lp_build_emit_fetch(bld_base, inst, 2, 0);
-       data1 = LLVMBuildBitCast(builder, tmp, bld_base->uint_bld.elem_type, "");
+       data1 = LLVMBuildBitCast(builder, tmp, ctx->i32, "");
 
        if (inst->Instruction.Opcode == TGSI_OPCODE_ATOMCAS) {
                tmp = lp_build_emit_fetch(bld_base, inst, 3, 0);
-               data2 = LLVMBuildBitCast(builder, tmp, bld_base->uint_bld.elem_type, "");
+               data2 = LLVMBuildBitCast(builder, tmp, ctx->i32, "");
        }
 
        /* llvm.amdgcn.image/buffer.atomic.cmpswap reflect the hardware order
@@ -3930,9 +3985,9 @@ static void atomic_fetch_args(
                rsrc = shader_buffer_fetch_rsrc(ctx, &inst->Src[0]);
 
                tmp = lp_build_emit_fetch(bld_base, inst, 1, 0);
-               offset = LLVMBuildBitCast(builder, tmp, bld_base->uint_bld.elem_type, "");
+               offset = LLVMBuildBitCast(builder, tmp, ctx->i32, "");
 
-               buffer_append_args(ctx, emit_data, rsrc, bld_base->uint_bld.zero,
+               buffer_append_args(ctx, emit_data, rsrc, ctx->i32_0,
                                   offset, true, false);
        } else if (inst->Src[0].Register.File == TGSI_FILE_IMAGE) {
                unsigned target = inst->Memory.Texture;
@@ -3943,7 +3998,7 @@ static void atomic_fetch_args(
 
                if (target == TGSI_TEXTURE_BUFFER) {
                        buffer_append_args(ctx, emit_data, rsrc, coords,
-                                          bld_base->uint_bld.zero, true, false);
+                                          ctx->i32_0, true, false);
                } else {
                        emit_data->args[emit_data->arg_count++] = coords;
                        emit_data->args[emit_data->arg_count++] = rsrc;
@@ -4028,7 +4083,7 @@ static void atomic_emit(
                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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        const struct tgsi_full_instruction * inst = emit_data->inst;
        char intrinsic_name[40];
@@ -4059,10 +4114,10 @@ static void atomic_emit(
        }
 
        tmp = lp_build_intrinsic(
-               builder, intrinsic_name, bld_base->uint_bld.elem_type,
+               builder, intrinsic_name, ctx->i32,
                emit_data->args, emit_data->arg_count, 0);
        emit_data->output[emit_data->chan] =
-               LLVMBuildBitCast(builder, tmp, bld_base->base.elem_type, "");
+               LLVMBuildBitCast(builder, tmp, ctx->f32, "");
 }
 
 static void set_tex_fetch_args(struct si_shader_context *ctx,
@@ -4096,6 +4151,35 @@ static void set_tex_fetch_args(struct si_shader_context *ctx,
        memcpy(emit_data->args, &args, sizeof(args));
 }
 
+static LLVMValueRef fix_resinfo(struct si_shader_context *ctx,
+                               unsigned target, LLVMValueRef out)
+{
+       LLVMBuilderRef builder = ctx->gallivm.builder;
+
+       /* 1D textures are allocated and used as 2D on GFX9. */
+        if (ctx->screen->b.chip_class >= GFX9 &&
+           (target == TGSI_TEXTURE_1D_ARRAY ||
+            target == TGSI_TEXTURE_SHADOW1D_ARRAY)) {
+               LLVMValueRef layers =
+                       LLVMBuildExtractElement(builder, out,
+                                               LLVMConstInt(ctx->i32, 2, 0), "");
+               out = LLVMBuildInsertElement(builder, out, layers,
+                                            ctx->i32_1, "");
+       }
+
+       /* Divide the number of layers by 6 to get the number of cubes. */
+       if (target == TGSI_TEXTURE_CUBE_ARRAY ||
+           target == TGSI_TEXTURE_SHADOWCUBE_ARRAY) {
+               LLVMValueRef imm2 = LLVMConstInt(ctx->i32, 2, 0);
+
+               LLVMValueRef z = LLVMBuildExtractElement(builder, out, imm2, "");
+               z = LLVMBuildSDiv(builder, z, LLVMConstInt(ctx->i32, 6, 0), "");
+
+               out = LLVMBuildInsertElement(builder, out, z, imm2, "");
+       }
+       return out;
+}
+
 static void resq_fetch_args(
                struct lp_build_tgsi_context * bld_base,
                struct lp_build_emit_data * emit_data)
@@ -4125,7 +4209,7 @@ static void resq_fetch_args(
                image_fetch_rsrc(bld_base, reg, false, inst->Memory.Texture,
                                 &res_ptr);
                set_tex_fetch_args(ctx, emit_data, image_target,
-                                  res_ptr, NULL, &bld_base->uint_bld.zero, 1,
+                                  res_ptr, NULL, &ctx->i32_0, 1,
                                   0xf);
        }
 }
@@ -4136,14 +4220,14 @@ static void resq_emit(
                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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        const struct tgsi_full_instruction *inst = emit_data->inst;
        LLVMValueRef out;
 
        if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
                out = LLVMBuildExtractElement(builder, emit_data->args[0],
-                                             lp_build_const_int32(gallivm, 2), "");
+                                             LLVMConstInt(ctx->i32, 2, 0), "");
        } else if (inst->Memory.Texture == TGSI_TEXTURE_BUFFER) {
                out = get_buffer_size(bld_base, emit_data->args[0]);
        } else {
@@ -4153,15 +4237,7 @@ static void resq_emit(
                args.opcode = ac_image_get_resinfo;
                out = ac_build_image_opcode(&ctx->ac, &args);
 
-               /* Divide the number of layers by 6 to get the number of cubes. */
-               if (inst->Memory.Texture == TGSI_TEXTURE_CUBE_ARRAY) {
-                       LLVMValueRef imm2 = lp_build_const_int32(gallivm, 2);
-                       LLVMValueRef imm6 = lp_build_const_int32(gallivm, 6);
-
-                       LLVMValueRef z = LLVMBuildExtractElement(builder, out, imm2, "");
-                       z = LLVMBuildSDiv(builder, z, imm6, "");
-                       out = LLVMBuildInsertElement(builder, out, z, imm2, "");
-               }
+               out = fix_resinfo(ctx, inst->Memory.Texture, out);
        }
 
        emit_data->output[emit_data->chan] = out;
@@ -4179,9 +4255,9 @@ enum desc_type {
 /**
  * Load an image view, fmask view. or sampler state descriptor.
  */
-static LLVMValueRef load_sampler_desc_custom(struct si_shader_context *ctx,
-                                            LLVMValueRef list, LLVMValueRef index,
-                                            enum desc_type type)
+static LLVMValueRef load_sampler_desc(struct si_shader_context *ctx,
+                                     LLVMValueRef list, LLVMValueRef index,
+                                     enum desc_type type)
 {
        struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
@@ -4194,14 +4270,14 @@ static LLVMValueRef load_sampler_desc_custom(struct si_shader_context *ctx,
        case DESC_BUFFER:
                /* The buffer is in [4:7]. */
                index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->i32, 4, 0), "");
-               index = LLVMBuildAdd(builder, index, LLVMConstInt(ctx->i32, 1, 0), "");
+               index = LLVMBuildAdd(builder, index, ctx->i32_1, "");
                list = LLVMBuildPointerCast(builder, list,
                                            const_array(ctx->v4i32, 0), "");
                break;
        case DESC_FMASK:
                /* The FMASK is at [8:15]. */
                index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->i32, 2, 0), "");
-               index = LLVMBuildAdd(builder, index, LLVMConstInt(ctx->i32, 1, 0), "");
+               index = LLVMBuildAdd(builder, index, ctx->i32_1, "");
                break;
        case DESC_SAMPLER:
                /* The sampler state is at [12:15]. */
@@ -4215,15 +4291,6 @@ static LLVMValueRef load_sampler_desc_custom(struct si_shader_context *ctx,
        return ac_build_indexed_load_const(&ctx->ac, list, index);
 }
 
-static LLVMValueRef load_sampler_desc(struct si_shader_context *ctx,
-                                    LLVMValueRef index, enum desc_type type)
-{
-       LLVMValueRef list = LLVMGetParam(ctx->main_fn,
-                                        SI_PARAM_SAMPLERS);
-
-       return load_sampler_desc_custom(ctx, list, index, type);
-}
-
 /* Disable anisotropic filtering if BASE_LEVEL == LAST_LEVEL.
  *
  * SI-CI:
@@ -4247,10 +4314,10 @@ static LLVMValueRef sici_fix_sampler_aniso(struct si_shader_context *ctx,
        img7 = LLVMBuildExtractElement(builder, res,
                                       LLVMConstInt(ctx->i32, 7, 0), "");
        samp0 = LLVMBuildExtractElement(builder, samp,
-                                       LLVMConstInt(ctx->i32, 0, 0), "");
+                                       ctx->i32_0, "");
        samp0 = LLVMBuildAnd(builder, samp0, img7, "");
        return LLVMBuildInsertElement(builder, samp, samp0,
-                                     LLVMConstInt(ctx->i32, 0, 0), "");
+                                     ctx->i32_0, "");
 }
 
 static void tex_fetch_ptrs(
@@ -4259,30 +4326,29 @@ static void tex_fetch_ptrs(
        LLVMValueRef *res_ptr, LLVMValueRef *samp_ptr, LLVMValueRef *fmask_ptr)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
+       LLVMValueRef list = LLVMGetParam(ctx->main_fn, SI_PARAM_SAMPLERS);
        const struct tgsi_full_instruction *inst = emit_data->inst;
+       const struct tgsi_full_src_register *reg;
        unsigned target = inst->Texture.Texture;
        unsigned sampler_src;
-       unsigned sampler_index;
        LLVMValueRef index;
 
        sampler_src = emit_data->inst->Instruction.NumSrcRegs - 1;
-       sampler_index = emit_data->inst->Src[sampler_src].Register.Index;
-
-       if (emit_data->inst->Src[sampler_src].Register.Indirect) {
-               const struct tgsi_full_src_register *reg = &emit_data->inst->Src[sampler_src];
+       reg = &emit_data->inst->Src[sampler_src];
 
+       if (reg->Register.Indirect) {
                index = get_bounded_indirect_index(ctx,
                                                   &reg->Indirect,
                                                   reg->Register.Index,
                                                   SI_NUM_SAMPLERS);
        } else {
-               index = LLVMConstInt(ctx->i32, sampler_index, 0);
+               index = LLVMConstInt(ctx->i32, reg->Register.Index, 0);
        }
 
        if (target == TGSI_TEXTURE_BUFFER)
-               *res_ptr = load_sampler_desc(ctx, index, DESC_BUFFER);
+               *res_ptr = load_sampler_desc(ctx, list, index, DESC_BUFFER);
        else
-               *res_ptr = load_sampler_desc(ctx, index, DESC_IMAGE);
+               *res_ptr = load_sampler_desc(ctx, list, index, DESC_IMAGE);
 
        if (samp_ptr)
                *samp_ptr = NULL;
@@ -4292,10 +4358,12 @@ static void tex_fetch_ptrs(
        if (target == TGSI_TEXTURE_2D_MSAA ||
            target == TGSI_TEXTURE_2D_ARRAY_MSAA) {
                if (fmask_ptr)
-                       *fmask_ptr = load_sampler_desc(ctx, index, DESC_FMASK);
+                       *fmask_ptr = load_sampler_desc(ctx, list, index,
+                                                      DESC_FMASK);
        } else if (target != TGSI_TEXTURE_BUFFER) {
                if (samp_ptr) {
-                       *samp_ptr = load_sampler_desc(ctx, index, DESC_SAMPLER);
+                       *samp_ptr = load_sampler_desc(ctx, list, index,
+                                                     DESC_SAMPLER);
                        *samp_ptr = sici_fix_sampler_aniso(ctx, *res_ptr, *samp_ptr);
                }
        }
@@ -4343,23 +4411,9 @@ static void txq_emit(const struct lp_build_tgsi_action *action,
        memcpy(&args, emit_data->args, sizeof(args)); /* ugly */
 
        args.opcode = ac_image_get_resinfo;
-       emit_data->output[emit_data->chan] =
-               ac_build_image_opcode(&ctx->ac, &args);
+       LLVMValueRef result = ac_build_image_opcode(&ctx->ac, &args);
 
-       /* Divide the number of layers by 6 to get the number of cubes. */
-       if (target == TGSI_TEXTURE_CUBE_ARRAY ||
-           target == TGSI_TEXTURE_SHADOWCUBE_ARRAY) {
-               LLVMBuilderRef builder = bld_base->base.gallivm->builder;
-               LLVMValueRef two = lp_build_const_int32(bld_base->base.gallivm, 2);
-               LLVMValueRef six = lp_build_const_int32(bld_base->base.gallivm, 6);
-
-               LLVMValueRef v4 = emit_data->output[emit_data->chan];
-               LLVMValueRef z = LLVMBuildExtractElement(builder, v4, two, "");
-               z = LLVMBuildSDiv(builder, z, six, "");
-
-               emit_data->output[emit_data->chan] =
-                       LLVMBuildInsertElement(builder, v4, z, two, "");
-       }
+       emit_data->output[emit_data->chan] = fix_resinfo(ctx, target, result);
 }
 
 static void tex_fetch_args(
@@ -4367,7 +4421,7 @@ static void tex_fetch_args(
        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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        const struct tgsi_full_instruction *inst = emit_data->inst;
        unsigned opcode = inst->Instruction.Opcode;
        unsigned target = inst->Texture.Texture;
@@ -4388,7 +4442,7 @@ static void tex_fetch_args(
                emit_data->dst_type = ctx->v4f32;
                emit_data->args[0] = LLVMBuildBitCast(gallivm->builder, res_ptr,
                                                      ctx->v16i8, "");
-               emit_data->args[1] = bld_base->uint_bld.zero;
+               emit_data->args[1] = ctx->i32_0;
                emit_data->args[2] = lp_build_emit_fetch(bld_base, emit_data->inst, 0, TGSI_CHAN_X);
                emit_data->arg_count = 3;
                return;
@@ -4411,7 +4465,9 @@ static void tex_fetch_args(
                coords[3] = bld_base->base.one;
 
        /* Pack offsets. */
-       if (has_offset && opcode != TGSI_OPCODE_TXF) {
+       if (has_offset &&
+           opcode != TGSI_OPCODE_TXF &&
+           opcode != TGSI_OPCODE_TXF_LZ) {
                /* The offsets are six-bit signed integers packed like this:
                 *   X=[5:0], Y=[13:8], and Z=[21:16].
                 */
@@ -4423,10 +4479,10 @@ static void tex_fetch_args(
                        offset[chan] = lp_build_emit_fetch_texoffset(bld_base,
                                                                     emit_data->inst, 0, chan);
                        offset[chan] = LLVMBuildAnd(gallivm->builder, offset[chan],
-                                                   lp_build_const_int32(gallivm, 0x3f), "");
+                                                   LLVMConstInt(ctx->i32, 0x3f, 0), "");
                        if (chan)
                                offset[chan] = LLVMBuildShl(gallivm->builder, offset[chan],
-                                                           lp_build_const_int32(gallivm, chan*8), "");
+                                                           LLVMConstInt(ctx->i32, chan*8, 0), "");
                }
 
                pack = LLVMBuildOr(gallivm->builder, offset[0], offset[1], "");
@@ -4466,11 +4522,12 @@ static void tex_fetch_args(
 
        /* Pack user derivatives */
        if (opcode == TGSI_OPCODE_TXD) {
-               int param, num_src_deriv_channels;
+               int param, num_src_deriv_channels, num_dst_deriv_channels;
 
                switch (target) {
                case TGSI_TEXTURE_3D:
                        num_src_deriv_channels = 3;
+                       num_dst_deriv_channels = 3;
                        num_deriv_channels = 3;
                        break;
                case TGSI_TEXTURE_2D:
@@ -4480,6 +4537,7 @@ static void tex_fetch_args(
                case TGSI_TEXTURE_2D_ARRAY:
                case TGSI_TEXTURE_SHADOW2D_ARRAY:
                        num_src_deriv_channels = 2;
+                       num_dst_deriv_channels = 2;
                        num_deriv_channels = 2;
                        break;
                case TGSI_TEXTURE_CUBE:
@@ -4488,6 +4546,7 @@ static void tex_fetch_args(
                case TGSI_TEXTURE_SHADOWCUBE_ARRAY:
                        /* Cube derivatives will be converted to 2D. */
                        num_src_deriv_channels = 3;
+                       num_dst_deriv_channels = 3;
                        num_deriv_channels = 2;
                        break;
                case TGSI_TEXTURE_1D:
@@ -4495,16 +4554,31 @@ static void tex_fetch_args(
                case TGSI_TEXTURE_1D_ARRAY:
                case TGSI_TEXTURE_SHADOW1D_ARRAY:
                        num_src_deriv_channels = 1;
-                       num_deriv_channels = 1;
+
+                       /* 1D textures are allocated and used as 2D on GFX9. */
+                       if (ctx->screen->b.chip_class >= GFX9) {
+                               num_dst_deriv_channels = 2;
+                               num_deriv_channels = 2;
+                       } else {
+                               num_dst_deriv_channels = 1;
+                               num_deriv_channels = 1;
+                       }
                        break;
                default:
                        unreachable("invalid target");
                }
 
-               for (param = 0; param < 2; param++)
+               for (param = 0; param < 2; param++) {
                        for (chan = 0; chan < num_src_deriv_channels; chan++)
-                               derivs[param * num_src_deriv_channels + chan] =
+                               derivs[param * num_dst_deriv_channels + chan] =
                                        lp_build_emit_fetch(bld_base, inst, param+1, chan);
+
+                       /* Fill in the rest with zeros. */
+                       for (chan = num_src_deriv_channels;
+                            chan < num_dst_deriv_channels; chan++)
+                               derivs[param * num_dst_deriv_channels + chan] =
+                                       bld_base->base.zero;
+               }
        }
 
        if (target == TGSI_TEXTURE_CUBE ||
@@ -4528,6 +4602,27 @@ static void tex_fetch_args(
        if (num_coords > 2)
                address[count++] = coords[2];
 
+       /* 1D textures are allocated and used as 2D on GFX9. */
+       if (ctx->screen->b.chip_class >= GFX9) {
+               LLVMValueRef filler;
+
+               /* Use 0.5, so that we don't sample the border color. */
+               if (opcode == TGSI_OPCODE_TXF)
+                       filler = ctx->i32_0;
+               else
+                       filler = LLVMConstReal(ctx->f32, 0.5);
+
+               if (target == TGSI_TEXTURE_1D ||
+                   target == TGSI_TEXTURE_SHADOW1D) {
+                       address[count++] = filler;
+               } else if (target == TGSI_TEXTURE_1D_ARRAY ||
+                          target == TGSI_TEXTURE_SHADOW1D_ARRAY) {
+                       address[count] = address[count - 1];
+                       address[count - 1] = filler;
+                       count++;
+               }
+       }
+
        /* Pack LOD or sample index */
        if (opcode == TGSI_OPCODE_TXL || opcode == TGSI_OPCODE_TXF)
                address[count++] = coords[3];
@@ -4560,7 +4655,6 @@ static void tex_fetch_args(
         */
        if (target == TGSI_TEXTURE_2D_MSAA ||
            target == TGSI_TEXTURE_2D_ARRAY_MSAA) {
-               struct lp_build_context *uint_bld = &bld_base->uint_bld;
                struct lp_build_emit_data txf_emit_data = *emit_data;
                LLVMValueRef txf_address[4];
                /* We only need .xy for non-arrays, and .xyz for arrays. */
@@ -4569,8 +4663,8 @@ static void tex_fetch_args(
 
                memcpy(txf_address, address, sizeof(txf_address));
 
-               /* Read FMASK using TXF. */
-               inst.Instruction.Opcode = TGSI_OPCODE_TXF;
+               /* Read FMASK using TXF_LZ. */
+               inst.Instruction.Opcode = TGSI_OPCODE_TXF_LZ;
                inst.Texture.Texture = target;
                txf_emit_data.inst = &inst;
                txf_emit_data.chan = 0;
@@ -4587,7 +4681,7 @@ static void tex_fetch_args(
                LLVMValueRef fmask =
                        LLVMBuildExtractElement(gallivm->builder,
                                                txf_emit_data.output[0],
-                                               uint_bld->zero, "");
+                                               ctx->i32_0, "");
 
                unsigned sample_chan = txf_count; /* the sample index is last */
 
@@ -4609,11 +4703,11 @@ static void tex_fetch_args(
 
                LLVMValueRef fmask_word1 =
                        LLVMBuildExtractElement(gallivm->builder, fmask_desc,
-                                               uint_bld->one, "");
+                                               ctx->i32_1, "");
 
                LLVMValueRef word1_is_nonzero =
                        LLVMBuildICmp(gallivm->builder, LLVMIntNE,
-                                     fmask_word1, uint_bld->zero, "");
+                                     fmask_word1, ctx->i32_0, "");
 
                /* Replace the MSAA sample index. */
                address[sample_chan] =
@@ -4621,7 +4715,8 @@ static void tex_fetch_args(
                                        final_sample, address[sample_chan], "");
        }
 
-       if (opcode == TGSI_OPCODE_TXF) {
+       if (opcode == TGSI_OPCODE_TXF ||
+           opcode == TGSI_OPCODE_TXF_LZ) {
                /* add tex offsets */
                if (inst->Texture.NumOffsets) {
                        struct lp_build_context *uint_bld = &bld_base->uint_bld;
@@ -4722,8 +4817,7 @@ static void si_lower_gather4_integer(struct si_shader_context *ctx,
                txq_emit_data.inst = &txq_inst;
                txq_emit_data.dst_type = ctx->v4i32;
                set_tex_fetch_args(ctx, &txq_emit_data, target,
-                                  args->resource, NULL,
-                                  &ctx->bld_base.uint_bld.zero,
+                                  args->resource, NULL, &ctx->i32_0,
                                   1, 0xf);
                txq_emit(NULL, &ctx->bld_base, &txq_emit_data);
 
@@ -4783,7 +4877,9 @@ static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
 
        switch (opcode) {
        case TGSI_OPCODE_TXF:
-               args.opcode = target == TGSI_TEXTURE_2D_MSAA ||
+       case TGSI_OPCODE_TXF_LZ:
+               args.opcode = opcode == TGSI_OPCODE_TXF_LZ ||
+                             target == TGSI_TEXTURE_2D_MSAA ||
                              target == TGSI_TEXTURE_2D_ARRAY_MSAA ?
                                      ac_image_load : ac_image_load_mip;
                args.compare = false;
@@ -4800,6 +4896,9 @@ static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
                if (ctx->type != PIPE_SHADER_FRAGMENT)
                        args.level_zero = true;
                break;
+       case TGSI_OPCODE_TEX_LZ:
+               args.level_zero = true;
+               break;
        case TGSI_OPCODE_TXB:
        case TGSI_OPCODE_TXB2:
                assert(ctx->type == PIPE_SHADER_FRAGMENT);
@@ -4822,7 +4921,8 @@ static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
        }
 
        /* The hardware needs special lowering for Gather4 with integer formats. */
-       if (opcode == TGSI_OPCODE_TG4) {
+       if (ctx->screen->b.chip_class <= VI &&
+           opcode == TGSI_OPCODE_TG4) {
                struct tgsi_shader_info *info = &ctx->shader->selector->info;
                /* This will also work with non-constant indexing because of how
                 * glsl_to_tgsi works and we intent to preserve that behavior.
@@ -4847,7 +4947,7 @@ static void si_llvm_emit_txqs(
        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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        LLVMValueRef res, samples;
        LLVMValueRef res_ptr, samp_ptr, fmask_ptr = NULL;
@@ -4859,12 +4959,12 @@ static void si_llvm_emit_txqs(
        res = LLVMBuildBitCast(builder, res_ptr, ctx->v8i32, "");
        samples = LLVMBuildExtractElement(
                builder, res,
-               lp_build_const_int32(gallivm, 3), "");
+               LLVMConstInt(ctx->i32, 3, 0), "");
        samples = LLVMBuildLShr(builder, samples,
-                               lp_build_const_int32(gallivm, 16), "");
+                               LLVMConstInt(ctx->i32, 16, 0), "");
        samples = LLVMBuildAnd(builder, samples,
-                              lp_build_const_int32(gallivm, 0xf), "");
-       samples = LLVMBuildShl(builder, lp_build_const_int32(gallivm, 1),
+                              LLVMConstInt(ctx->i32, 0xf, 0), "");
+       samples = LLVMBuildShl(builder, ctx->i32_1,
                               samples, "");
 
        emit_data->output[emit_data->chan] = samples;
@@ -4876,7 +4976,7 @@ static void si_llvm_emit_ddxy(
        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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        unsigned opcode = emit_data->info->opcode;
        LLVMValueRef val;
        int idx;
@@ -4908,7 +5008,7 @@ static LLVMValueRef si_llvm_emit_ddxy_interp(
        LLVMValueRef interp_ij)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef result[4], a;
        unsigned i;
 
@@ -4927,7 +5027,7 @@ static void interp_fetch_args(
        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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        const struct tgsi_full_instruction *inst = emit_data->inst;
 
        if (inst->Instruction.Opcode == TGSI_OPCODE_INTERP_OFFSET) {
@@ -4942,7 +5042,7 @@ static void interp_fetch_args(
        } else if (inst->Instruction.Opcode == TGSI_OPCODE_INTERP_SAMPLE) {
                LLVMValueRef sample_position;
                LLVMValueRef sample_id;
-               LLVMValueRef halfval = lp_build_const_float(gallivm, 0.5f);
+               LLVMValueRef halfval = LLVMConstReal(ctx->f32, 0.5f);
 
                /* fetch sample ID, then fetch its sample position,
                 * and place into first two channels.
@@ -4955,12 +5055,12 @@ static void interp_fetch_args(
 
                emit_data->args[0] = LLVMBuildExtractElement(gallivm->builder,
                                                             sample_position,
-                                                            lp_build_const_int32(gallivm, 0), "");
+                                                            ctx->i32_0, "");
 
                emit_data->args[0] = LLVMBuildFSub(gallivm->builder, emit_data->args[0], halfval, "");
                emit_data->args[1] = LLVMBuildExtractElement(gallivm->builder,
                                                             sample_position,
-                                                            lp_build_const_int32(gallivm, 1), "");
+                                                            ctx->i32_1, "");
                emit_data->args[1] = LLVMBuildFSub(gallivm->builder, emit_data->args[1], halfval, "");
                emit_data->arg_count = 2;
        }
@@ -4972,8 +5072,7 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct si_shader *shader = ctx->shader;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       struct lp_build_context *uint = &bld_base->uint_bld;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMValueRef interp_param;
        const struct tgsi_full_instruction *inst = emit_data->inst;
        int input_index = inst->Src[0].Register.Index;
@@ -5001,7 +5100,7 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
        else
                interp_param = NULL;
 
-       attr_number = lp_build_const_int32(gallivm, input_index);
+       attr_number = LLVMConstInt(ctx->i32, input_index, 0);
 
        if (inst->Instruction.Opcode == TGSI_OPCODE_INTERP_OFFSET ||
            inst->Instruction.Opcode == TGSI_OPCODE_INTERP_SAMPLE) {
@@ -5017,8 +5116,8 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
                 * interp_param.J = ddy * offset/sample.y + temp1;
                 */
                for (i = 0; i < 2; i++) {
-                       LLVMValueRef ix_ll = lp_build_const_int32(gallivm, i);
-                       LLVMValueRef iy_ll = lp_build_const_int32(gallivm, i + 2);
+                       LLVMValueRef ix_ll = LLVMConstInt(ctx->i32, i, 0);
+                       LLVMValueRef iy_ll = LLVMConstInt(ctx->i32, i + 2, 0);
                        LLVMValueRef ddx_el = LLVMBuildExtractElement(gallivm->builder,
                                                                      ddxy_out, ix_ll, "");
                        LLVMValueRef ddy_el = LLVMBuildExtractElement(gallivm->builder,
@@ -5038,7 +5137,7 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
 
                        ij_out[i] = LLVMBuildFAdd(gallivm->builder, temp2, temp1, "");
                }
-               interp_param = lp_build_gather_values(bld_base->base.gallivm, ij_out, 2);
+               interp_param = lp_build_gather_values(gallivm, ij_out, 2);
        }
 
        for (chan = 0; chan < 4; chan++) {
@@ -5046,26 +5145,164 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
                unsigned schan;
 
                schan = tgsi_util_get_full_src_register_swizzle(&inst->Src[0], chan);
-               llvm_chan = lp_build_const_int32(gallivm, schan);
+               llvm_chan = LLVMConstInt(ctx->i32, schan, 0);
 
                if (interp_param) {
                        interp_param = LLVMBuildBitCast(gallivm->builder,
                                interp_param, LLVMVectorType(ctx->f32, 2), "");
                        LLVMValueRef i = LLVMBuildExtractElement(
-                               gallivm->builder, interp_param, uint->zero, "");
+                               gallivm->builder, interp_param, ctx->i32_0, "");
                        LLVMValueRef j = LLVMBuildExtractElement(
-                               gallivm->builder, interp_param, uint->one, "");
+                               gallivm->builder, interp_param, ctx->i32_1, "");
                        emit_data->output[chan] = ac_build_fs_interp(&ctx->ac,
                                llvm_chan, attr_number, params,
                                i, j);
                } else {
                        emit_data->output[chan] = ac_build_fs_interp_mov(&ctx->ac,
-                               lp_build_const_int32(gallivm, 2), /* P0 */
+                               LLVMConstInt(ctx->i32, 2, 0), /* P0 */
                                llvm_chan, attr_number, params);
                }
        }
 }
 
+static LLVMValueRef si_emit_ballot(struct si_shader_context *ctx,
+                                  LLVMValueRef value)
+{
+       struct gallivm_state *gallivm = &ctx->gallivm;
+       LLVMValueRef args[3] = {
+               value,
+               ctx->i32_0,
+               LLVMConstInt(ctx->i32, LLVMIntNE, 0)
+       };
+
+       /* We currently have no other way to prevent LLVM from lifting the icmp
+        * calls to a dominating basic block.
+        */
+       emit_optimization_barrier(ctx, &args[0]);
+
+       if (LLVMTypeOf(args[0]) != ctx->i32)
+               args[0] = LLVMBuildBitCast(gallivm->builder, args[0], ctx->i32, "");
+
+       return lp_build_intrinsic(gallivm->builder,
+                                 "llvm.amdgcn.icmp.i32",
+                                 ctx->i64, args, 3,
+                                 LP_FUNC_ATTR_NOUNWIND |
+                                 LP_FUNC_ATTR_READNONE |
+                                 LP_FUNC_ATTR_CONVERGENT);
+}
+
+static void vote_all_emit(
+       const struct lp_build_tgsi_action *action,
+       struct lp_build_tgsi_context *bld_base,
+       struct lp_build_emit_data *emit_data)
+{
+       struct si_shader_context *ctx = si_shader_context(bld_base);
+       struct gallivm_state *gallivm = &ctx->gallivm;
+       LLVMValueRef active_set, vote_set;
+       LLVMValueRef tmp;
+
+       active_set = si_emit_ballot(ctx, ctx->i32_1);
+       vote_set = si_emit_ballot(ctx, emit_data->args[0]);
+
+       tmp = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, vote_set, active_set, "");
+       emit_data->output[emit_data->chan] =
+               LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, "");
+}
+
+static void vote_any_emit(
+       const struct lp_build_tgsi_action *action,
+       struct lp_build_tgsi_context *bld_base,
+       struct lp_build_emit_data *emit_data)
+{
+       struct si_shader_context *ctx = si_shader_context(bld_base);
+       struct gallivm_state *gallivm = &ctx->gallivm;
+       LLVMValueRef vote_set;
+       LLVMValueRef tmp;
+
+       vote_set = si_emit_ballot(ctx, emit_data->args[0]);
+
+       tmp = LLVMBuildICmp(gallivm->builder, LLVMIntNE,
+                           vote_set, LLVMConstInt(ctx->i64, 0, 0), "");
+       emit_data->output[emit_data->chan] =
+               LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, "");
+}
+
+static void vote_eq_emit(
+       const struct lp_build_tgsi_action *action,
+       struct lp_build_tgsi_context *bld_base,
+       struct lp_build_emit_data *emit_data)
+{
+       struct si_shader_context *ctx = si_shader_context(bld_base);
+       struct gallivm_state *gallivm = &ctx->gallivm;
+       LLVMValueRef active_set, vote_set;
+       LLVMValueRef all, none, tmp;
+
+       active_set = si_emit_ballot(ctx, ctx->i32_1);
+       vote_set = si_emit_ballot(ctx, emit_data->args[0]);
+
+       all = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, vote_set, active_set, "");
+       none = LLVMBuildICmp(gallivm->builder, LLVMIntEQ,
+                            vote_set, LLVMConstInt(ctx->i64, 0, 0), "");
+       tmp = LLVMBuildOr(gallivm->builder, all, none, "");
+       emit_data->output[emit_data->chan] =
+               LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, "");
+}
+
+static void ballot_emit(
+       const struct lp_build_tgsi_action *action,
+       struct lp_build_tgsi_context *bld_base,
+       struct lp_build_emit_data *emit_data)
+{
+       struct si_shader_context *ctx = si_shader_context(bld_base);
+       LLVMBuilderRef builder = ctx->gallivm.builder;
+       LLVMValueRef tmp;
+
+       tmp = lp_build_emit_fetch(bld_base, emit_data->inst, 0, TGSI_CHAN_X);
+       tmp = si_emit_ballot(ctx, tmp);
+       tmp = LLVMBuildBitCast(builder, tmp, ctx->v2i32, "");
+
+       emit_data->output[0] = LLVMBuildExtractElement(builder, tmp, ctx->i32_0, "");
+       emit_data->output[1] = LLVMBuildExtractElement(builder, tmp, ctx->i32_1, "");
+}
+
+static void read_invoc_fetch_args(
+       struct lp_build_tgsi_context *bld_base,
+       struct lp_build_emit_data *emit_data)
+{
+       emit_data->args[0] = lp_build_emit_fetch(bld_base, emit_data->inst,
+                                                0, emit_data->src_chan);
+
+       /* Always read the source invocation (= lane) from the X channel. */
+       emit_data->args[1] = lp_build_emit_fetch(bld_base, emit_data->inst,
+                                                1, TGSI_CHAN_X);
+       emit_data->arg_count = 2;
+}
+
+static void read_lane_emit(
+       const struct lp_build_tgsi_action *action,
+       struct lp_build_tgsi_context *bld_base,
+       struct lp_build_emit_data *emit_data)
+{
+       struct si_shader_context *ctx = si_shader_context(bld_base);
+       LLVMBuilderRef builder = ctx->gallivm.builder;
+
+       /* We currently have no other way to prevent LLVM from lifting the icmp
+        * calls to a dominating basic block.
+        */
+       emit_optimization_barrier(ctx, &emit_data->args[0]);
+
+       for (unsigned i = 0; i < emit_data->arg_count; ++i) {
+               emit_data->args[i] = LLVMBuildBitCast(builder, emit_data->args[i],
+                                                     ctx->i32, "");
+       }
+
+       emit_data->output[emit_data->chan] =
+               ac_build_intrinsic(&ctx->ac, action->intr_name,
+                                  ctx->i32, emit_data->args, emit_data->arg_count,
+                                  AC_FUNC_ATTR_READNONE |
+                                  AC_FUNC_ATTR_CONVERGENT);
+}
+
 static unsigned si_llvm_get_stream(struct lp_build_tgsi_context *bld_base,
                                       struct lp_build_emit_data *emit_data)
 {
@@ -5091,7 +5328,7 @@ static void si_llvm_emit_vertex(
        struct lp_build_context *uint = &bld_base->uint_bld;
        struct si_shader *shader = ctx->shader;
        struct tgsi_shader_info *info = &shader->selector->info;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        struct lp_build_if_state if_state;
        LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
                                            SI_PARAM_GS2VS_OFFSET);
@@ -5117,14 +5354,14 @@ static void si_llvm_emit_vertex(
         * altogether.
         */
        can_emit = LLVMBuildICmp(gallivm->builder, LLVMIntULT, gs_next_vertex,
-                                lp_build_const_int32(gallivm,
-                                                     shader->selector->gs_max_out_vertices), "");
+                                LLVMConstInt(ctx->i32,
+                                             shader->selector->gs_max_out_vertices, 0), "");
 
        bool use_kill = !info->writes_memory;
        if (use_kill) {
                kill = lp_build_select(&bld_base->base, can_emit,
-                                      lp_build_const_float(gallivm, 1.0f),
-                                      lp_build_const_float(gallivm, -1.0f));
+                                      LLVMConstReal(ctx->f32, 1.0f),
+                                      LLVMConstReal(ctx->f32, -1.0f));
 
                ac_build_kill(&ctx->ac, kill);
        } else {
@@ -5142,8 +5379,8 @@ static void si_llvm_emit_vertex(
 
                        LLVMValueRef out_val = LLVMBuildLoad(gallivm->builder, out_ptr[chan], "");
                        LLVMValueRef voffset =
-                               lp_build_const_int32(gallivm, offset *
-                                                    shader->selector->gs_max_out_vertices);
+                               LLVMConstInt(ctx->i32, offset *
+                                            shader->selector->gs_max_out_vertices, 0);
                        offset++;
 
                        voffset = lp_build_add(uint, voffset, gs_next_vertex);
@@ -5160,7 +5397,7 @@ static void si_llvm_emit_vertex(
        }
 
        gs_next_vertex = lp_build_add(uint, gs_next_vertex,
-                                     lp_build_const_int32(gallivm, 1));
+                                     ctx->i32_1);
 
        LLVMBuildStore(gallivm->builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
 
@@ -5191,7 +5428,7 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
                                 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;
+       struct gallivm_state *gallivm = &ctx->gallivm;
 
        /* SI only (thanks to a hw bug workaround):
         * The real barrier instruction isn’t needed, because an entire patch
@@ -5246,7 +5483,7 @@ static void si_create_function(struct si_shader_context *ctx,
                if (LLVMGetTypeKind(LLVMTypeOf(P)) == LLVMPointerTypeKind) {
                        lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_BYVAL);
                        lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_NOALIAS);
-                       lp_add_attr_dereferenceable(P, UINT64_MAX);
+                       ac_add_attr_dereferenceable(P, UINT64_MAX);
                } else
                        lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_INREG);
        }
@@ -5323,11 +5560,9 @@ 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->bld_base;
-       struct lp_build_context *uint = &bld_base->uint_bld;
 
        unsigned lds_size = ctx->screen->b.chip_class >= CIK ? 65536 : 32768;
-       ctx->lds = LLVMBuildIntToPtr(gallivm->builder, uint->zero,
+       ctx->lds = LLVMBuildIntToPtr(gallivm->builder, ctx->i32_0,
                LLVMPointerType(LLVMArrayType(ctx->i32, lds_size / 4), LOCAL_ADDR_SPACE),
                "tess_lds");
 }
@@ -5352,7 +5587,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->bld_base;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        struct si_shader *shader = ctx->shader;
        LLVMTypeRef params[SI_NUM_PARAMS + SI_MAX_ATTRIBS], v3i32;
        LLVMTypeRef returns[16+32*4];
@@ -5617,7 +5852,7 @@ static void create_function(struct si_shader_context *ctx)
  */
 static void preload_ring_buffers(struct si_shader_context *ctx)
 {
-       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
 
        LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
@@ -5631,21 +5866,20 @@ static void preload_ring_buffers(struct si_shader_context *ctx)
                unsigned ring =
                        ctx->type == PIPE_SHADER_GEOMETRY ? SI_GS_RING_ESGS
                                                             : SI_ES_RING_ESGS;
-               LLVMValueRef offset = lp_build_const_int32(gallivm, ring);
+               LLVMValueRef offset = LLVMConstInt(ctx->i32, ring, 0);
 
                ctx->esgs_ring =
                        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_RING_GSVS);
+               LLVMValueRef offset = LLVMConstInt(ctx->i32, SI_RING_GSVS, 0);
 
                ctx->gsvs_ring[0] =
                        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 offset = LLVMConstInt(ctx->i32, SI_RING_GSVS, 0);
                LLVMValueRef base_ring;
 
                base_ring = ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
@@ -5679,20 +5913,20 @@ static void preload_ring_buffers(struct si_shader_context *ctx)
                        num_records = 64;
 
                        ring = LLVMBuildBitCast(builder, base_ring, v2i64, "");
-                       tmp = LLVMBuildExtractElement(builder, ring, uint->zero, "");
+                       tmp = LLVMBuildExtractElement(builder, ring, ctx->i32_0, "");
                        tmp = LLVMBuildAdd(builder, tmp,
                                           LLVMConstInt(ctx->i64,
                                                        stream_offset, 0), "");
                        stream_offset += stride * 64;
 
-                       ring = LLVMBuildInsertElement(builder, ring, tmp, uint->zero, "");
+                       ring = LLVMBuildInsertElement(builder, ring, tmp, ctx->i32_0, "");
                        ring = LLVMBuildBitCast(builder, ring, ctx->v4i32, "");
-                       tmp = LLVMBuildExtractElement(builder, ring, uint->one, "");
+                       tmp = LLVMBuildExtractElement(builder, ring, ctx->i32_1, "");
                        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, tmp, ctx->i32_1, "");
                        ring = LLVMBuildInsertElement(builder, ring,
                                        LLVMConstInt(ctx->i32, num_records, 0),
                                        LLVMConstInt(ctx->i32, 2, 0), "");
@@ -5720,8 +5954,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->bld_base;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
+       struct gallivm_state *gallivm = &ctx->gallivm;
        LLVMBuilderRef builder = gallivm->builder;
        LLVMValueRef slot, desc, offset, row, bit, address[2];
 
@@ -5733,7 +5966,7 @@ static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx,
        address[1] = unpack_param(ctx, param_pos_fixed_pt, 16, 5);
 
        /* Load the buffer descriptor. */
-       slot = lp_build_const_int32(gallivm, SI_PS_CONST_POLY_STIPPLE);
+       slot = LLVMConstInt(ctx->i32, SI_PS_CONST_POLY_STIPPLE, 0);
        desc = ac_build_indexed_load_const(&ctx->ac, param_rw_buffers, slot);
 
        /* The stipple pattern is 32x32, each row has 32 bits. */
@@ -5891,6 +6124,12 @@ int si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader)
        assert((!prolog && !epilog) || !mainb->rodata_size);
        assert(!epilog || !epilog->rodata_size);
 
+       /* GFX9 can fetch at most 128 bytes past the end of the shader.
+        * Prevent VM faults.
+        */
+       if (sscreen->b.chip_class >= GFX9)
+               bo_size += 128;
+
        r600_resource_reference(&shader->bo, NULL);
        shader->bo = (struct r600_resource*)
                     pipe_buffer_create(&sscreen->b.b, 0,
@@ -6255,7 +6494,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
        if (gs_selector->so.num_outputs)
                stream_id = unpack_param(&ctx, ctx.param_streamout_config, 24, 2);
        else
-               stream_id = uint->zero;
+               stream_id = ctx.i32_0;
 
        /* Fill in output information. */
        for (i = 0; i < gsinfo->num_outputs; ++i) {
@@ -6285,7 +6524,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
                        continue;
 
                bb = LLVMInsertBasicBlockInContext(gallivm->context, end_bb, "out");
-               LLVMAddCase(switch_inst, lp_build_const_int32(gallivm, stream), bb);
+               LLVMAddCase(switch_inst, LLVMConstInt(ctx.i32, stream, 0), bb);
                LLVMPositionBuilderAtEnd(builder, bb);
 
                /* Fetch vertex data from GSVS ring */
@@ -6305,7 +6544,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
                                outputs[i].values[chan] =
                                        ac_build_buffer_load(&ctx.ac,
                                                             ctx.gsvs_ring[0], 1,
-                                                            uint->zero, voffset,
+                                                            ctx.i32_0, voffset,
                                                             soffset, 0, 1, 1, true);
                        }
                }
@@ -6330,14 +6569,14 @@ si_generate_gs_copy_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, PIPE_SHADER_GEOMETRY))
-               ac_dump_module(bld_base->base.gallivm->module);
+               ac_dump_module(ctx.gallivm.module);
 
        si_llvm_finalize_module(&ctx,
                r600_extra_shader_checks(&sscreen->b, PIPE_SHADER_GEOMETRY));
 
        r = si_compile_llvm(sscreen, &ctx.shader->binary,
                            &ctx.shader->config, ctx.tm,
-                           bld_base->base.gallivm->module,
+                           ctx.gallivm.module,
                            debug, PIPE_SHADER_GEOMETRY,
                            "GS Copy Shader");
        if (!r) {
@@ -6454,11 +6693,13 @@ static void si_init_shader_ctx(struct si_shader_context *ctx,
        bld_base->op_actions[TGSI_OPCODE_INTERP_OFFSET] = interp_action;
 
        bld_base->op_actions[TGSI_OPCODE_TEX] = tex_action;
+       bld_base->op_actions[TGSI_OPCODE_TEX_LZ] = tex_action;
        bld_base->op_actions[TGSI_OPCODE_TEX2] = tex_action;
        bld_base->op_actions[TGSI_OPCODE_TXB] = tex_action;
        bld_base->op_actions[TGSI_OPCODE_TXB2] = tex_action;
        bld_base->op_actions[TGSI_OPCODE_TXD] = tex_action;
        bld_base->op_actions[TGSI_OPCODE_TXF] = tex_action;
+       bld_base->op_actions[TGSI_OPCODE_TXF_LZ] = tex_action;
        bld_base->op_actions[TGSI_OPCODE_TXL] = tex_action;
        bld_base->op_actions[TGSI_OPCODE_TXL2] = tex_action;
        bld_base->op_actions[TGSI_OPCODE_TXP] = tex_action;
@@ -6500,16 +6741,31 @@ static void si_init_shader_ctx(struct si_shader_context *ctx,
 
        bld_base->op_actions[TGSI_OPCODE_MEMBAR].emit = membar_emit;
 
+       bld_base->op_actions[TGSI_OPCODE_CLOCK].emit = clock_emit;
+
        bld_base->op_actions[TGSI_OPCODE_DDX].emit = si_llvm_emit_ddxy;
        bld_base->op_actions[TGSI_OPCODE_DDY].emit = si_llvm_emit_ddxy;
        bld_base->op_actions[TGSI_OPCODE_DDX_FINE].emit = si_llvm_emit_ddxy;
        bld_base->op_actions[TGSI_OPCODE_DDY_FINE].emit = si_llvm_emit_ddxy;
 
+       bld_base->op_actions[TGSI_OPCODE_VOTE_ALL].emit = vote_all_emit;
+       bld_base->op_actions[TGSI_OPCODE_VOTE_ANY].emit = vote_any_emit;
+       bld_base->op_actions[TGSI_OPCODE_VOTE_EQ].emit = vote_eq_emit;
+       bld_base->op_actions[TGSI_OPCODE_BALLOT].emit = ballot_emit;
+       bld_base->op_actions[TGSI_OPCODE_READ_FIRST].intr_name = "llvm.amdgcn.readfirstlane";
+       bld_base->op_actions[TGSI_OPCODE_READ_FIRST].emit = read_lane_emit;
+       bld_base->op_actions[TGSI_OPCODE_READ_INVOC].intr_name = "llvm.amdgcn.readlane";
+       bld_base->op_actions[TGSI_OPCODE_READ_INVOC].fetch_args = read_invoc_fetch_args;
+       bld_base->op_actions[TGSI_OPCODE_READ_INVOC].emit = read_lane_emit;
+
        bld_base->op_actions[TGSI_OPCODE_EMIT].emit = si_llvm_emit_vertex;
        bld_base->op_actions[TGSI_OPCODE_ENDPRIM].emit = si_llvm_emit_primitive;
        bld_base->op_actions[TGSI_OPCODE_BARRIER].emit = si_llvm_emit_barrier;
 }
 
+#define EXP_TARGET (HAVE_LLVM >= 0x0500 ? 0 : 3)
+#define EXP_OUT0 (HAVE_LLVM >= 0x0500 ? 2 : 5)
+
 /* Return true if the PARAM export has been eliminated. */
 static bool si_eliminate_const_output(struct si_shader_context *ctx,
                                      LLVMValueRef inst, unsigned offset)
@@ -6521,7 +6777,7 @@ static bool si_eliminate_const_output(struct si_shader_context *ctx,
 
        for (i = 0; i < 4; i++) {
                LLVMBool loses_info;
-               LLVMValueRef p = LLVMGetOperand(inst, (HAVE_LLVM >= 0x0500 ? 2 : 5) + i);
+               LLVMValueRef p = LLVMGetOperand(inst, EXP_OUT0 + i);
 
                /* It's a constant expression. Undef outputs are eliminated too. */
                if (LLVMIsUndef(p)) {
@@ -6607,10 +6863,10 @@ static void si_eliminate_const_vs_outputs(struct si_shader_context *ctx)
                        /* Check if this is an export instruction. */
                        if ((num_args != 9 && num_args != 8) ||
                            (strcmp(name, "llvm.SI.export") &&
-                            strcmp(name, "llvm.amdgcn.exp.")))
+                            strcmp(name, "llvm.amdgcn.exp.f32")))
                                continue;
 
-                       LLVMValueRef arg = LLVMGetOperand(cur, HAVE_LLVM >= 0x0500 ? 0 : 3);
+                       LLVMValueRef arg = LLVMGetOperand(cur, EXP_TARGET);
                        unsigned target = LLVMConstIntGetZExtValue(arg);
 
                        if (target < V_008DFC_SQ_EXP_PARAM)
@@ -6652,7 +6908,7 @@ static void si_eliminate_const_vs_outputs(struct si_shader_context *ctx)
                                if (current_offset[out] != offset)
                                        continue;
 
-                               LLVMSetOperand(exports.inst[i], 3,
+                               LLVMSetOperand(exports.inst[i], EXP_TARGET,
                                               LLVMConstInt(ctx->i32,
                                                            V_008DFC_SQ_EXP_PARAM + new_count, 0));
                                shader->info.vs_output_param_offset[out] = new_count;
@@ -6742,7 +6998,7 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx,
                int i;
                for (i = 0; i < 4; i++) {
                        ctx->gs_next_vertex[i] =
-                               lp_build_alloca(bld_base->base.gallivm,
+                               lp_build_alloca(&ctx->gallivm,
                                                ctx->i32, "");
                }
        }
@@ -7222,7 +7478,6 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 {
        struct si_shader_selector *sel = shader->selector;
        struct si_shader_context ctx;
-       struct lp_build_tgsi_context *bld_base;
        LLVMModuleRef mod;
        int r = -1;
 
@@ -7242,7 +7497,6 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 
        shader->info.uses_instanceid = sel->info.uses_instanceid;
 
-       bld_base = &ctx.bld_base;
        ctx.load_system_value = declare_system_value;
 
        if (!si_compile_tgsi_main(&ctx, shader)) {
@@ -7335,7 +7589,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
                si_build_wrapper_function(&ctx, parts, need_prolog ? 3 : 2, need_prolog ? 1 : 0);
        }
 
-       mod = bld_base->base.gallivm->module;
+       mod = ctx.gallivm.module;
 
        /* Dump LLVM IR before any optimization passes */
        if (sscreen->b.debug_flags & DBG_PREOPT_IR &&
@@ -7472,7 +7726,7 @@ si_get_shader_part(struct si_screen *sscreen,
        /* Find existing. */
        for (result = *list; result; result = result->next) {
                if (memcmp(&result->key, key, sizeof(*key)) == 0) {
-                       pipe_mutex_unlock(sscreen->shader_parts_mutex);
+                       mtx_unlock(&sscreen->shader_parts_mutex);
                        return result;
                }
        }
@@ -7526,7 +7780,7 @@ si_get_shader_part(struct si_screen *sscreen,
 
 out:
        si_llvm_dispose(&ctx);
-       pipe_mutex_unlock(sscreen->shader_parts_mutex);
+       mtx_unlock(&sscreen->shader_parts_mutex);
        return result;
 }