radeonsi: change the bit-packing of LS out/TCS in data
[mesa.git] / src / gallium / drivers / radeonsi / si_shader.c
index 055c8f7daed52e3c2735ef84989d98236620d680..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;
@@ -352,13 +351,12 @@ static void declare_input_vs(
        LLVMValueRef t_offset;
        LLVMValueRef t_list;
        LLVMValueRef vertex_index;
-       LLVMValueRef args[3];
        LLVMValueRef input[3];
 
        /* Load the T list */
        t_list_ptr = LLVMGetParam(ctx->main_fn, SI_PARAM_VERTEX_BUFFERS);
 
-       t_offset = lp_build_const_int32(gallivm, input_index);
+       t_offset = LLVMConstInt(ctx->i32, input_index, 0);
 
        t_list = ac_build_indexed_load_const(&ctx->ac, t_list_ptr, t_offset);
 
@@ -393,21 +391,17 @@ static void declare_input_vs(
                fetch_stride = 0;
        }
 
-       args[0] = t_list;
-       args[2] = vertex_index;
-
        for (unsigned i = 0; i < num_fetches; i++) {
-               args[1] = LLVMConstInt(ctx->i32, fetch_stride * i, 0);
+               LLVMValueRef voffset = LLVMConstInt(ctx->i32, fetch_stride * i, 0);
 
-               input[i] = lp_build_intrinsic(gallivm->builder,
-                       "llvm.SI.vs.load.input", ctx->v4f32, args, 3,
-                       LP_FUNC_ATTR_READNONE |
-                       LP_FUNC_ATTR_LEGACY);
+               input[i] = ac_build_buffer_load_format(&ctx->ac, t_list,
+                                                      vertex_index, voffset,
+                                                      true);
        }
 
        /* 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, "");
        }
@@ -552,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:
@@ -569,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;
        }
 }
 
@@ -581,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;
 }
 
@@ -605,7 +599,7 @@ static LLVMValueRef get_bounded_indirect_index(struct si_shader_context *ctx,
         * - SI & CIK hang
         * - VI crashes
         */
-       if (HAVE_LLVM <= 0x0308)
+       if (HAVE_LLVM == 0x0308)
                return LLVMGetUndef(ctx->i32);
 
        return si_llvm_bound_index(ctx, result, num);
@@ -621,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;
@@ -649,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,
@@ -684,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 {
@@ -694,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
@@ -720,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;
 
@@ -729,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, "");
@@ -764,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;
@@ -780,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. */
@@ -809,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),
@@ -826,35 +819,35 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
 static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
                                 enum tgsi_opcode_type type, unsigned swizzle,
                                 LLVMValueRef buffer, LLVMValueRef offset,
-                                LLVMValueRef 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);
 
        if (swizzle == ~0) {
                value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
-                                            0, 1, 0);
+                                            0, 1, 0, readonly_memory);
 
                return LLVMBuildBitCast(gallivm->builder, value, vec_type, "");
        }
 
        if (!tgsi_type_is_64bit(type)) {
                value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
-                                            0, 1, 0);
+                                            0, 1, 0, readonly_memory);
 
                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,
-                                 swizzle * 4, 1, 0);
+                                 swizzle * 4, 1, 0, readonly_memory);
 
        value2 = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
-                                  swizzle * 4 + 4, 1, 0);
+                                  swizzle * 4 + 4, 1, 0, readonly_memory);
 
        return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
 }
@@ -871,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) {
@@ -880,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);
        }
@@ -912,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,
@@ -930,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);
 
@@ -963,18 +956,17 @@ 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);
 
-       return buffer_load(bld_base, type, swizzle, buffer, base, addr);
+       return buffer_load(bld_base, type, swizzle, buffer, base, addr, true);
 }
 
 static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
@@ -983,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;
@@ -1027,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);
@@ -1037,7 +1029,7 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
                LLVMValueRef value = dst[chan_index];
 
                if (inst->Instruction.Saturate)
-                       value = ac_emit_clamp(&ctx->ac, value);
+                       value = ac_build_clamp(&ctx->ac, value);
 
                /* Skip LDS stores if there is no LDS read of this output. */
                if (!skip_lds_store)
@@ -1047,17 +1039,17 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
                values[chan_index] = value;
 
                if (inst->Dst[0].Register.WriteMask != 0xF && !is_tess_factor) {
-                       ac_build_tbuffer_store_dwords(&ctx->ac, buffer, value, 1,
-                                                     buf_addr, base,
-                                                     4 * chan_index);
+                       ac_build_buffer_store_dword(&ctx->ac, buffer, value, 1,
+                                                   buf_addr, base,
+                                                   4 * chan_index, 1, 0, true, false);
                }
        }
 
        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_tbuffer_store_dwords(&ctx->ac, buffer, value, 4, buf_addr,
-                                             base, 0);
+               ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buf_addr,
+                                           base, 0, 1, 0, true, false);
        }
 }
 
@@ -1067,13 +1059,11 @@ 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;
-       LLVMValueRef vtx_offset;
-       LLVMValueRef args[9];
+       struct gallivm_state *gallivm = &ctx->gallivm;
+       LLVMValueRef vtx_offset, soffset;
        unsigned vtx_offset_param;
        struct tgsi_shader_info *info = &shader->selector->info;
        unsigned semantic_name = info->input_semantic_name[reg->Register.Index];
@@ -1093,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);
        }
 
@@ -1111,29 +1101,17 @@ static LLVMValueRef fetch_input_gs(
                                      4);
 
        param = si_shader_io_get_unique_index(semantic_name, semantic_index);
-       args[0] = ctx->esgs_ring;
-       args[1] = vtx_offset;
-       args[2] = lp_build_const_int32(gallivm, (param * 4 + swizzle) * 256);
-       args[3] = uint->zero;
-       args[4] = uint->one;  /* OFFEN */
-       args[5] = uint->zero; /* IDXEN */
-       args[6] = uint->one;  /* GLC */
-       args[7] = uint->zero; /* SLC */
-       args[8] = uint->zero; /* TFE */
-
-       value = lp_build_intrinsic(gallivm->builder,
-                                  "llvm.SI.buffer.load.dword.i32.i32",
-                                  ctx->i32, args, 9,
-                                  LP_FUNC_ATTR_READONLY |
-                                  LP_FUNC_ATTR_LEGACY);
+       soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle) * 256, 0);
+
+       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;
-               args[2] = lp_build_const_int32(gallivm, (param * 4 + swizzle + 1) * 256);
-               value2 = lp_build_intrinsic(gallivm->builder,
-                                           "llvm.SI.buffer.load.dword.i32.i32",
-                                           ctx->i32, args, 9,
-                                           LP_FUNC_ATTR_READONLY |
-                                           LP_FUNC_ATTR_LEGACY);
+               soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle + 1) * 256, 0);
+
+               value2 = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1,
+                                             ctx->i32_0, vtx_offset, soffset,
+                                             0, 1, 0, true);
                return si_llvm_emit_fetch_64bit(bld_base, type,
                                                value, value2);
        }
@@ -1196,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;
 
@@ -1220,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 &&
@@ -1244,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) {
@@ -1262,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);
                        }
 
@@ -1277,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);
                        }
                }
@@ -1304,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;
 
@@ -1353,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);
 }
 
 
@@ -1375,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;
 
@@ -1449,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");
@@ -1458,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);
@@ -1470,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;
@@ -1496,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
                };
@@ -1536,14 +1506,14 @@ 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,
-                                   ~0, buffer, base, addr);
+               value = buffer_load(&ctx->bld_base, TGSI_TYPE_FLOAT,
+                                   ~0, buffer, base, addr, true);
 
                break;
        }
@@ -1554,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:
@@ -1588,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:
@@ -1619,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;
@@ -1680,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;
@@ -1701,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);
        }
@@ -1724,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 */
@@ -1753,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;
@@ -1822,33 +1829,29 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
                        };
                        LLVMValueRef packed;
 
-                       packed = lp_build_intrinsic(base->gallivm->builder,
-                                                   "llvm.SI.packf16",
-                                                   ctx->i32, pack_args, 2,
-                                                   LP_FUNC_ATTR_READNONE |
-                                                   LP_FUNC_ATTR_LEGACY);
+                       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;
 
        case V_028714_SPI_SHADER_UNORM16_ABGR:
                for (chan = 0; chan < 4; chan++) {
-                       val[chan] = ac_emit_clamp(&ctx->ac, values[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:
@@ -1856,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++) {
@@ -1896,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++) {
@@ -1925,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;
        }
 
@@ -1941,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,
@@ -1954,14 +1956,12 @@ 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));
 
-               lp_build_intrinsic(gallivm->builder, "llvm.AMDGPU.kill",
-                                  ctx->voidt, &arg, 1, LP_FUNC_ATTR_LEGACY);
+               ac_build_kill(&ctx->ac, arg);
        } else {
-               lp_build_intrinsic(gallivm->builder, "llvm.AMDGPU.kilp",
-                                  ctx->voidt, NULL, 0, LP_FUNC_ATTR_LEGACY);
+               ac_build_kill(&ctx->ac, NULL);
        }
 }
 
@@ -1970,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 */
@@ -1986,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, "");
@@ -2002,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 ++) {
@@ -2012,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++) {
@@ -2102,11 +2102,11 @@ static void emit_streamout_output(struct si_shader_context *ctx,
                break;
        }
 
-       ac_build_tbuffer_store_dwords(&ctx->ac, so_buffers[buf_idx],
-                                     vdata, num_comps,
-                                     so_write_offsets[buf_idx],
-                                     LLVMConstInt(ctx->i32, 0, 0),
-                                     stream_out->dst_offset * 4);
+       ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf_idx],
+                                   vdata, num_comps,
+                                   so_write_offsets[buf_idx],
+                                   ctx->i32_0,
+                                   stream_out->dst_offset * 4, 1, 1, true, false);
 }
 
 /**
@@ -2163,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);
 
@@ -2300,7 +2300,7 @@ handle_semantic:
                        memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS],
                               &args, sizeof(args));
                } else {
-                       ac_emit_export(&ctx->ac, &args);
+                       ac_build_export(&ctx->ac, &args);
                }
 
                if (semantic_name == TGSI_SEMANTIC_CLIPDIST) {
@@ -2348,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, "");
                }
@@ -2384,7 +2384,7 @@ handle_semantic:
                        /* Specify that this is the last export */
                        pos_args[i].done = 1;
 
-               ac_emit_export(&ctx->ac, &pos_args[i]);
+               ac_build_export(&ctx->ac, &pos_args[i]);
        }
 }
 
@@ -2395,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;
@@ -2404,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);
@@ -2419,19 +2419,19 @@ 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);
 
-               ac_build_tbuffer_store_dwords(&ctx->ac, buffer, value, 4, buffer_addr,
-                                             buffer_offset, 0);
+               ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buffer_addr,
+                                           buffer_offset, 0, 1, 0, true, false);
        }
 }
 
@@ -2441,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;
@@ -2459,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) {
@@ -2491,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);
@@ -2530,31 +2530,34 @@ 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_tbuffer_store_dwords(&ctx->ac, buffer,
-                                     lp_build_const_int32(gallivm, 0x80000000),
-                                     1, lp_build_const_int32(gallivm, 0), tf_base, 0);
+       ac_build_buffer_store_dword(&ctx->ac, buffer,
+                                   LLVMConstInt(ctx->i32, 0x80000000, 0),
+                                   1, ctx->i32_0, tf_base,
+                                   0, 1, 0, true, false);
 
        lp_build_endif(&inner_if_ctx);
 
        /* Store the tessellation factors. */
-       ac_build_tbuffer_store_dwords(&ctx->ac, buffer, vec0,
-                                     MIN2(stride, 4), byteoffset, tf_base, 4);
+       ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
+                                   MIN2(stride, 4), byteoffset, tf_base,
+                                   4, 1, 0, true, false);
        if (vec1)
-               ac_build_tbuffer_store_dwords(&ctx->ac, buffer, vec1,
-                                             stride - 4, byteoffset, tf_base, 20);
+               ac_build_buffer_store_dword(&ctx->ac, buffer, vec1,
+                                           stride - 4, byteoffset, tf_base,
+                                           20, 1, 0, true, false);
 
        /* Store the tess factors into the offchip buffer if TES reads them. */
        if (shader->key.part.tcs.epilog.tes_reads_tess_factors) {
@@ -2574,9 +2577,9 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
                outer_vec = lp_build_gather_values(gallivm, outer,
                                                   util_next_power_of_two(outer_comps));
 
-               ac_build_tbuffer_store_dwords(&ctx->ac, buf, outer_vec,
-                                             outer_comps, tf_outer_offset,
-                                             base, 0);
+               ac_build_buffer_store_dword(&ctx->ac, buf, outer_vec,
+                                           outer_comps, tf_outer_offset,
+                                           base, 0, 1, 0, true, false);
                if (inner_comps) {
                        param_inner = si_shader_io_get_unique_index(
                                              TGSI_SEMANTIC_TESSINNER, 0);
@@ -2585,9 +2588,9 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
 
                        inner_vec = inner_comps == 1 ? inner[0] :
                                    lp_build_gather_values(gallivm, inner, inner_comps);
-                       ac_build_tbuffer_store_dwords(&ctx->ac, buf, inner_vec,
-                                                     inner_comps, tf_inner_offset,
-                                                     base, 0);
+                       ac_build_buffer_store_dword(&ctx->ac, buf, inner_vec,
+                                                   inner_comps, tf_inner_offset,
+                                                   base, 0, 1, 0, true, false);
                }
        }
 
@@ -2608,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;
@@ -2619,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, "");
 
@@ -2655,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, "");
 
@@ -2672,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,
@@ -2684,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,
@@ -2707,14 +2710,11 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base)
                        LLVMValueRef out_val = LLVMBuildLoad(gallivm->builder, out_ptr[chan], "");
                        out_val = LLVMBuildBitCast(gallivm->builder, out_val, ctx->i32, "");
 
-                       ac_build_tbuffer_store(&ctx->ac,
-                                              ctx->esgs_ring,
-                                              out_val, 1,
-                                              LLVMGetUndef(ctx->i32), soffset,
-                                              (4 * param_index + chan) * 4,
-                                              V_008F0C_BUF_DATA_FORMAT_32,
-                                              V_008F0C_BUF_NUM_FORMAT_UINT,
-                                              0, 0, 1, 1, 0);
+                       ac_build_buffer_store_dword(&ctx->ac,
+                                                   ctx->esgs_ring,
+                                                   out_val, 1, NULL, soffset,
+                                                   (4 * param_index + chan) * 4,
+                                                   1, 1, true, true);
                }
        }
 }
@@ -2723,14 +2723,14 @@ static void si_llvm_emit_gs_epilogue(struct lp_build_tgsi_context *bld_base)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
 
-       ac_emit_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_NOP | AC_SENDMSG_GS_DONE,
-                       LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
+       ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_NOP | AC_SENDMSG_GS_DONE,
+                        LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
 }
 
 static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
 {
        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;
@@ -2768,7 +2768,7 @@ static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
                        for (j = 0; j < 4; j++) {
                                addr = ctx->outputs[i][j];
                                val = LLVMBuildLoad(gallivm->builder, addr, "");
-                               val = ac_emit_clamp(&ctx->ac, val);
+                               val = ac_build_clamp(&ctx->ac, val);
                                LLVMBuildStore(gallivm->builder, val, addr);
                        }
                }
@@ -2863,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;
@@ -2913,7 +2913,7 @@ static void si_export_mrt_color(struct lp_build_tgsi_context *bld_base,
        /* Clamp color */
        if (ctx->shader->key.part.ps.epilog.clamp_color)
                for (i = 0; i < 4; i++)
-                       color[i] = ac_emit_clamp(&ctx->ac, color[i]);
+                       color[i] = ac_build_clamp(&ctx->ac, color[i]);
 
        /* Alpha to one */
        if (ctx->shader->key.part.ps.epilog.alpha_to_one)
@@ -2972,7 +2972,7 @@ static void si_emit_ps_exports(struct si_shader_context *ctx,
                               struct si_ps_exports *exp)
 {
        for (unsigned i = 0; i < exp->num; i++)
-               ac_emit_export(&ctx->ac, &exp->args[i]);
+               ac_build_export(&ctx->ac, &exp->args[i]);
 }
 
 static void si_export_null(struct lp_build_tgsi_context *bld_base)
@@ -2991,7 +2991,7 @@ static void si_export_null(struct lp_build_tgsi_context *bld_base)
        args.out[2] = base->undef; /* B */
        args.out[3] = base->undef; /* A */
 
-       ac_emit_export(&ctx->ac, &args);
+       ac_build_export(&ctx->ac, &args);
 }
 
 /**
@@ -3011,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] = {};
@@ -3097,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, "");
        }
@@ -3129,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
@@ -3150,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);
@@ -3181,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)
@@ -3254,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.
  */
@@ -3267,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);
 
@@ -3279,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:
@@ -3297,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(
@@ -3321,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);
@@ -3331,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];
 
@@ -3423,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;
@@ -3438,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;
@@ -3450,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);
@@ -3462,8 +3522,25 @@ static void load_fetch_args(
        }
 }
 
+static unsigned get_load_intr_attribs(bool readonly_memory)
+{
+       /* READNONE means writes can't affect it, while READONLY means that
+        * writes can affect it. */
+       return readonly_memory && HAVE_LLVM >= 0x0400 ?
+                                LP_FUNC_ATTR_READNONE :
+                                LP_FUNC_ATTR_READONLY;
+}
+
+static unsigned get_store_intr_attribs(bool writeonly_memory)
+{
+       return writeonly_memory && HAVE_LLVM >= 0x0400 ?
+                                 LP_FUNC_ATTR_INACCESSIBLE_MEM_ONLY :
+                                 LP_FUNC_ATTR_WRITEONLY;
+}
+
 static void load_emit_buffer(struct si_shader_context *ctx,
-                            struct lp_build_emit_data *emit_data)
+                            struct lp_build_emit_data *emit_data,
+                            bool readonly_memory)
 {
        const struct tgsi_full_instruction *inst = emit_data->inst;
        struct gallivm_state *gallivm = &ctx->gallivm;
@@ -3491,7 +3568,7 @@ static void load_emit_buffer(struct si_shader_context *ctx,
        emit_data->output[emit_data->chan] = lp_build_intrinsic(
                        builder, intrinsic_name, dst_type,
                        emit_data->args, emit_data->arg_count,
-                       LP_FUNC_ATTR_READONLY);
+                       get_load_intr_attribs(readonly_memory));
 }
 
 static LLVMValueRef get_memory_ptr(struct si_shader_context *ctx,
@@ -3519,52 +3596,87 @@ 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];
+/**
+ * Return true if the memory accessed by a LOAD or STORE instruction is
+ * read-only or write-only, respectively.
+ *
+ * \param shader_buffers_reverse_access_mask
+ *     For LOAD, set this to (store | atomic) slot usage in the shader.
+ *     For STORE, set this to (load | atomic) slot usage in the shader.
+ * \param images_reverse_access_mask  Same as above, but for images.
+ */
+static bool is_oneway_access_only(const struct tgsi_full_instruction *inst,
+                                 const struct tgsi_shader_info *info,
+                                 unsigned shader_buffers_reverse_access_mask,
+                                 unsigned images_reverse_access_mask)
+{
+       /* RESTRICT means NOALIAS.
+        * If there are no writes, we can assume the accessed memory is read-only.
+        * If there are no reads, we can assume the accessed memory is write-only.
+        */
+       if (inst->Memory.Qualifier & TGSI_MEMORY_RESTRICT) {
+               unsigned reverse_access_mask;
+
+               if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
+                       reverse_access_mask = shader_buffers_reverse_access_mask;
+               } else if (inst->Memory.Texture == TGSI_TEXTURE_BUFFER) {
+                       reverse_access_mask = info->images_buffers &
+                                             images_reverse_access_mask;
+               } else {
+                       reverse_access_mask = ~info->images_buffers &
+                                             images_reverse_access_mask;
+               }
 
-       ac_build_type_name_for_intr(coords_type, coords_type_name,
-                           sizeof(coords_type_name));
+               if (inst->Src[0].Register.Indirect) {
+                       if (!reverse_access_mask)
+                               return true;
+               } else {
+                       if (!(reverse_access_mask &
+                             (1u << inst->Src[0].Register.Index)))
+                               return true;
+               }
+       }
 
-       if (HAVE_LLVM <= 0x0309) {
-               snprintf(out_name, out_len, "%s.%s", base_name, coords_type_name);
+       /* If there are no buffer writes (for both shader buffers & image
+        * buffers), it implies that buffer memory is read-only.
+        * If there are no buffer reads (for both shader buffers & image
+        * buffers), it implies that buffer memory is write-only.
+        *
+        * Same for the case when there are no writes/reads for non-buffer
+        * images.
+        */
+       if (inst->Src[0].Register.File == TGSI_FILE_BUFFER ||
+           (inst->Src[0].Register.File == TGSI_FILE_IMAGE &&
+            inst->Memory.Texture == TGSI_TEXTURE_BUFFER)) {
+               if (!shader_buffers_reverse_access_mask &&
+                   !(info->images_buffers & images_reverse_access_mask))
+                       return true;
        } 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);
+               if (!(~info->images_buffers & images_reverse_access_mask))
+                       return true;
        }
+       return false;
 }
 
 static void load_emit(
@@ -3573,10 +3685,12 @@ 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;
        char intrinsic_name[64];
+       bool readonly_memory = false;
 
        if (inst->Src[0].Register.File == TGSI_FILE_MEMORY) {
                load_emit_memory(ctx, emit_data);
@@ -3586,8 +3700,15 @@ static void load_emit(
        if (inst->Memory.Qualifier & TGSI_MEMORY_VOLATILE)
                emit_waitcnt(ctx, VM_CNT);
 
+       readonly_memory = !(inst->Memory.Qualifier & TGSI_MEMORY_VOLATILE) &&
+                         is_oneway_access_only(inst, info,
+                                               info->shader_buffers_store |
+                                               info->shader_buffers_atomic,
+                                               info->images_store |
+                                               info->images_atomic);
+
        if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
-               load_emit_buffer(ctx, emit_data);
+               load_emit_buffer(ctx, emit_data, readonly_memory);
                return;
        }
 
@@ -3596,19 +3717,19 @@ static void load_emit(
                        lp_build_intrinsic(
                                builder, "llvm.amdgcn.buffer.load.format.v4f32", emit_data->dst_type,
                                emit_data->args, emit_data->arg_count,
-                               LP_FUNC_ATTR_READONLY);
+                               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(
                                builder, intrinsic_name, emit_data->dst_type,
                                emit_data->args, emit_data->arg_count,
-                               LP_FUNC_ATTR_READONLY);
+                               get_load_intr_attribs(readonly_memory));
        }
 }
 
@@ -3617,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;
@@ -3644,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;
@@ -3665,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);
@@ -3680,12 +3801,12 @@ static void store_fetch_args(
 
 static void store_emit_buffer(
                struct si_shader_context *ctx,
-               struct lp_build_emit_data *emit_data)
+               struct lp_build_emit_data *emit_data,
+               bool writeonly_memory)
 {
        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;
@@ -3714,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";
                }
 
@@ -3738,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;
@@ -3746,7 +3867,8 @@ static void store_emit_buffer(
 
                lp_build_intrinsic(
                        builder, intrinsic_name, emit_data->dst_type,
-                       emit_data->args, emit_data->arg_count, 0);
+                       emit_data->args, emit_data->arg_count,
+                       get_store_intr_attribs(writeonly_memory));
        }
 }
 
@@ -3756,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);
        }
@@ -3781,11 +3902,13 @@ 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;
        unsigned target = inst->Memory.Texture;
        char intrinsic_name[64];
+       bool writeonly_memory = false;
 
        if (inst->Dst[0].Register.File == TGSI_FILE_MEMORY) {
                store_emit_memory(ctx, emit_data);
@@ -3795,8 +3918,14 @@ static void store_emit(
        if (inst->Memory.Qualifier & TGSI_MEMORY_VOLATILE)
                emit_waitcnt(ctx, VM_CNT);
 
+       writeonly_memory = is_oneway_access_only(inst, info,
+                                                info->shader_buffers_load |
+                                                info->shader_buffers_atomic,
+                                                info->images_load |
+                                                info->images_atomic);
+
        if (inst->Dst[0].Register.File == TGSI_FILE_BUFFER) {
-               store_emit_buffer(ctx, emit_data);
+               store_emit_buffer(ctx, emit_data, writeonly_memory);
                return;
        }
 
@@ -3804,18 +3933,20 @@ static void store_emit(
                emit_data->output[emit_data->chan] = lp_build_intrinsic(
                        builder, "llvm.amdgcn.buffer.store.format.v4f32",
                        emit_data->dst_type, emit_data->args,
-                       emit_data->arg_count, 0);
+                       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(
                                builder, intrinsic_name, emit_data->dst_type,
-                               emit_data->args, emit_data->arg_count, 0);
+                               emit_data->args, emit_data->arg_count,
+                               get_store_intr_attribs(writeonly_memory));
        }
 }
 
@@ -3824,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
@@ -3854,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;
@@ -3867,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;
@@ -3952,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];
@@ -3983,10 +4114,70 @@ 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,
+                              struct lp_build_emit_data *emit_data,
+                              unsigned target,
+                              LLVMValueRef res_ptr, LLVMValueRef samp_ptr,
+                              LLVMValueRef *param, unsigned count,
+                              unsigned dmask)
+{
+       struct gallivm_state *gallivm = &ctx->gallivm;
+       struct ac_image_args args = {};
+
+       /* Pad to power of two vector */
+       while (count < util_next_power_of_two(count))
+               param[count++] = LLVMGetUndef(ctx->i32);
+
+       if (count > 1)
+               args.addr = lp_build_gather_values(gallivm, param, count);
+       else
+               args.addr = param[0];
+
+       args.resource = res_ptr;
+       args.sampler = samp_ptr;
+       args.dmask = dmask;
+       args.unorm = target == TGSI_TEXTURE_RECT ||
+                    target == TGSI_TEXTURE_SHADOWRECT;
+       args.da = tgsi_is_array_sampler(target);
+
+       /* Ugly, but we seem to have no other choice right now. */
+       STATIC_ASSERT(sizeof(args) <= sizeof(emit_data->args));
+       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(
@@ -3994,7 +4185,6 @@ static void resq_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;
        const struct tgsi_full_instruction *inst = emit_data->inst;
        const struct tgsi_full_src_register *reg = &inst->Src[0];
 
@@ -4008,19 +4198,19 @@ static void resq_fetch_args(
                                 &emit_data->args[0]);
                emit_data->arg_count = 1;
        } else {
-               emit_data->args[0] = bld_base->uint_bld.zero; /* mip level */
+               LLVMValueRef res_ptr;
+               unsigned image_target;
+
+               if (inst->Memory.Texture == TGSI_TEXTURE_3D)
+                       image_target = TGSI_TEXTURE_2D_ARRAY;
+               else
+                       image_target = inst->Memory.Texture;
+
                image_fetch_rsrc(bld_base, reg, false, inst->Memory.Texture,
-                                &emit_data->args[1]);
-               emit_data->args[2] = lp_build_const_int32(gallivm, 15); /* dmask */
-               emit_data->args[3] = bld_base->uint_bld.zero; /* unorm */
-               emit_data->args[4] = bld_base->uint_bld.zero; /* r128 */
-               emit_data->args[5] = tgsi_is_array_image(inst->Memory.Texture) ?
-                       bld_base->uint_bld.one : bld_base->uint_bld.zero; /* da */
-               emit_data->args[6] = bld_base->uint_bld.zero; /* glc */
-               emit_data->args[7] = bld_base->uint_bld.zero; /* slc */
-               emit_data->args[8] = bld_base->uint_bld.zero; /* tfe */
-               emit_data->args[9] = bld_base->uint_bld.zero; /* lwe */
-               emit_data->arg_count = 10;
+                                &res_ptr);
+               set_tex_fetch_args(ctx, emit_data, image_target,
+                                  res_ptr, NULL, &ctx->i32_0, 1,
+                                  0xf);
        }
 }
 
@@ -4029,83 +4219,30 @@ static void resq_emit(
                struct lp_build_tgsi_context *bld_base,
                struct lp_build_emit_data *emit_data)
 {
-       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;
        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 {
-               out = lp_build_intrinsic(
-                       builder, "llvm.SI.getresinfo.i32", emit_data->dst_type,
-                       emit_data->args, emit_data->arg_count,
-                       LP_FUNC_ATTR_READNONE | LP_FUNC_ATTR_LEGACY);
+               struct ac_image_args 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);
+               memcpy(&args, emit_data->args, sizeof(args)); /* ugly */
+               args.opcode = ac_image_get_resinfo;
+               out = ac_build_image_opcode(&ctx->ac, &args);
 
-                       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;
 }
 
-static void set_tex_fetch_args(struct si_shader_context *ctx,
-                              struct lp_build_emit_data *emit_data,
-                              unsigned opcode, unsigned target,
-                              LLVMValueRef res_ptr, LLVMValueRef samp_ptr,
-                              LLVMValueRef *param, unsigned count,
-                              unsigned dmask)
-{
-       struct gallivm_state *gallivm = &ctx->gallivm;
-       unsigned num_args;
-       unsigned is_rect = target == TGSI_TEXTURE_RECT ||
-                          target == TGSI_TEXTURE_SHADOWRECT;
-
-       /* Pad to power of two vector */
-       while (count < util_next_power_of_two(count))
-               param[count++] = LLVMGetUndef(ctx->i32);
-
-       /* Texture coordinates. */
-       if (count > 1)
-               emit_data->args[0] = lp_build_gather_values(gallivm, param, count);
-       else
-               emit_data->args[0] = param[0];
-
-       /* Resource. */
-       emit_data->args[1] = res_ptr;
-       num_args = 2;
-
-       if (opcode == TGSI_OPCODE_TXF || opcode == TGSI_OPCODE_TXQ)
-               emit_data->dst_type = ctx->v4i32;
-       else {
-               emit_data->dst_type = ctx->v4f32;
-
-               emit_data->args[num_args++] = samp_ptr;
-       }
-
-       emit_data->args[num_args++] = lp_build_const_int32(gallivm, dmask);
-       emit_data->args[num_args++] = lp_build_const_int32(gallivm, is_rect); /* unorm */
-       emit_data->args[num_args++] = lp_build_const_int32(gallivm, 0); /* r128 */
-       emit_data->args[num_args++] = lp_build_const_int32(gallivm,
-                                       tgsi_is_array_sampler(target)); /* da */
-       emit_data->args[num_args++] = lp_build_const_int32(gallivm, 0); /* glc */
-       emit_data->args[num_args++] = lp_build_const_int32(gallivm, 0); /* slc */
-       emit_data->args[num_args++] = lp_build_const_int32(gallivm, 0); /* tfe */
-       emit_data->args[num_args++] = lp_build_const_int32(gallivm, 0); /* lwe */
-
-       emit_data->arg_count = num_args;
-}
-
 static const struct lp_build_tgsi_action tex_action;
 
 enum desc_type {
@@ -4118,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;
@@ -4133,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]. */
@@ -4154,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:
@@ -4186,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(
@@ -4198,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;
@@ -4231,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);
                }
        }
@@ -4261,7 +4390,7 @@ static void txq_fetch_args(
        /* Textures - set the mip level. */
        address = lp_build_emit_fetch(bld_base, inst, 0, TGSI_CHAN_X);
 
-       set_tex_fetch_args(ctx, emit_data, TGSI_OPCODE_TXQ, target, res_ptr,
+       set_tex_fetch_args(ctx, emit_data, target, res_ptr,
                           NULL, &address, 1, 0xf);
 }
 
@@ -4269,7 +4398,8 @@ static void txq_emit(const struct lp_build_tgsi_action *action,
                     struct lp_build_tgsi_context *bld_base,
                     struct lp_build_emit_data *emit_data)
 {
-       struct lp_build_context *base = &bld_base->base;
+       struct si_shader_context *ctx = si_shader_context(bld_base);
+       struct ac_image_args args;
        unsigned target = emit_data->inst->Texture.Texture;
 
        if (target == TGSI_TEXTURE_BUFFER) {
@@ -4278,25 +4408,12 @@ static void txq_emit(const struct lp_build_tgsi_action *action,
                return;
        }
 
-       emit_data->output[emit_data->chan] = lp_build_intrinsic(
-               base->gallivm->builder, "llvm.SI.getresinfo.i32",
-               emit_data->dst_type, emit_data->args, emit_data->arg_count,
-               LP_FUNC_ATTR_READNONE | LP_FUNC_ATTR_LEGACY);
-
-       /* 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);
+       memcpy(&args, emit_data->args, sizeof(args)); /* ugly */
 
-               LLVMValueRef v4 = emit_data->output[emit_data->chan];
-               LLVMValueRef z = LLVMBuildExtractElement(builder, v4, two, "");
-               z = LLVMBuildSDiv(builder, z, six, "");
+       args.opcode = ac_image_get_resinfo;
+       LLVMValueRef result = ac_build_image_opcode(&ctx->ac, &args);
 
-               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(
@@ -4304,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;
@@ -4325,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;
@@ -4348,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].
                 */
@@ -4360,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], "");
@@ -4396,18 +4515,19 @@ static void tex_fetch_args(
                 * Z32_FLOAT, but we don't know that here.
                 */
                if (ctx->screen->b.chip_class == VI)
-                       z = ac_emit_clamp(&ctx->ac, z);
+                       z = ac_build_clamp(&ctx->ac, z);
 
                address[count++] = z;
        }
 
        /* 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:
@@ -4417,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:
@@ -4425,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:
@@ -4432,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 ||
@@ -4465,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];
@@ -4497,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. */
@@ -4506,12 +4663,12 @@ 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;
-               set_tex_fetch_args(ctx, &txf_emit_data, TGSI_OPCODE_TXF,
+               set_tex_fetch_args(ctx, &txf_emit_data,
                                   target, fmask_ptr, NULL,
                                   txf_address, txf_count, 0xf);
                build_tex_intrinsic(&tex_action, bld_base, &txf_emit_data);
@@ -4524,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 */
 
@@ -4546,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] =
@@ -4558,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;
@@ -4620,7 +4778,7 @@ static void tex_fetch_args(
                dmask = 1 << gather_comp;
        }
 
-       set_tex_fetch_args(ctx, emit_data, opcode, target, res_ptr,
+       set_tex_fetch_args(ctx, emit_data, target, res_ptr,
                           samp_ptr, address, count, dmask);
 }
 
@@ -4634,30 +4792,32 @@ static void tex_fetch_args(
  * or (0.5 / size) from the normalized coordinates.
  */
 static void si_lower_gather4_integer(struct si_shader_context *ctx,
-                                    struct lp_build_emit_data *emit_data,
-                                    const char *intr_name,
-                                    unsigned coord_vgpr_index)
+                                    struct ac_image_args *args,
+                                    unsigned target)
 {
        LLVMBuilderRef builder = ctx->gallivm.builder;
-       LLVMValueRef coord = emit_data->args[0];
+       LLVMValueRef coord = args->addr;
        LLVMValueRef half_texel[2];
+       /* Texture coordinates start after:
+        *   {offset, bias, z-compare, derivatives}
+        * Only the offset and z-compare can occur here.
+        */
+       unsigned coord_vgpr_index = (int)args->offset + (int)args->compare;
        int c;
 
-       if (emit_data->inst->Texture.Texture == TGSI_TEXTURE_RECT ||
-           emit_data->inst->Texture.Texture == TGSI_TEXTURE_SHADOWRECT) {
+       if (target == TGSI_TEXTURE_RECT ||
+           target == TGSI_TEXTURE_SHADOWRECT) {
                half_texel[0] = half_texel[1] = LLVMConstReal(ctx->f32, -0.5);
        } else {
                struct tgsi_full_instruction txq_inst = {};
                struct lp_build_emit_data txq_emit_data = {};
 
                /* Query the texture size. */
-               txq_inst.Texture.Texture = emit_data->inst->Texture.Texture;
+               txq_inst.Texture.Texture = target;
                txq_emit_data.inst = &txq_inst;
                txq_emit_data.dst_type = ctx->v4i32;
-               set_tex_fetch_args(ctx, &txq_emit_data, TGSI_OPCODE_TXQ,
-                                  txq_inst.Texture.Texture,
-                                  emit_data->args[1], NULL,
-                                  &ctx->bld_base.uint_bld.zero,
+               set_tex_fetch_args(ctx, &txq_emit_data, target,
+                                  args->resource, NULL, &ctx->i32_0,
                                   1, 0xf);
                txq_emit(NULL, &ctx->bld_base, &txq_emit_data);
 
@@ -4686,11 +4846,7 @@ static void si_lower_gather4_integer(struct si_shader_context *ctx,
                coord = LLVMBuildInsertElement(builder, coord, tmp, index, "");
        }
 
-       emit_data->args[0] = coord;
-       emit_data->output[emit_data->chan] =
-               lp_build_intrinsic(builder, intr_name, emit_data->dst_type,
-                                  emit_data->args, emit_data->arg_count,
-                                  LP_FUNC_ATTR_READNONE | LP_FUNC_ATTR_LEGACY);
+       args->addr = coord;
 }
 
 static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
@@ -4698,75 +4854,75 @@ static void build_tex_intrinsic(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 lp_build_context *base = &bld_base->base;
        const struct tgsi_full_instruction *inst = emit_data->inst;
+       struct ac_image_args args;
        unsigned opcode = inst->Instruction.Opcode;
        unsigned target = inst->Texture.Texture;
-       char intr_name[127];
-       bool has_offset = inst->Texture.NumOffsets > 0;
-       bool is_shadow = tgsi_is_shadow_target(target);
-       char type[64];
-       const char *name = "llvm.SI.image.sample";
-       const char *infix = "";
 
        if (target == TGSI_TEXTURE_BUFFER) {
-               emit_data->output[emit_data->chan] = lp_build_intrinsic(
-                       base->gallivm->builder,
-                       "llvm.SI.vs.load.input", emit_data->dst_type,
-                       emit_data->args, emit_data->arg_count,
-                       LP_FUNC_ATTR_READNONE | LP_FUNC_ATTR_LEGACY);
+               emit_data->output[emit_data->chan] =
+                       ac_build_buffer_load_format(&ctx->ac,
+                                                   emit_data->args[0],
+                                                   emit_data->args[2],
+                                                   emit_data->args[1],
+                                                   true);
                return;
        }
 
+       memcpy(&args, emit_data->args, sizeof(args)); /* ugly */
+
+       args.opcode = ac_image_sample;
+       args.compare = tgsi_is_shadow_target(target);
+       args.offset = inst->Texture.NumOffsets > 0;
+
        switch (opcode) {
        case TGSI_OPCODE_TXF:
-               name = target == TGSI_TEXTURE_2D_MSAA ||
-                      target == TGSI_TEXTURE_2D_ARRAY_MSAA ?
-                              "llvm.SI.image.load" :
-                              "llvm.SI.image.load.mip";
-               is_shadow = false;
-               has_offset = false;
+       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;
+               args.offset = false;
                break;
        case TGSI_OPCODE_LODQ:
-               name = "llvm.SI.getlod";
-               is_shadow = false;
-               has_offset = false;
+               args.opcode = ac_image_get_lod;
+               args.compare = false;
+               args.offset = false;
                break;
        case TGSI_OPCODE_TEX:
        case TGSI_OPCODE_TEX2:
        case TGSI_OPCODE_TXP:
                if (ctx->type != PIPE_SHADER_FRAGMENT)
-                       infix = ".lz";
+                       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);
-               infix = ".b";
+               args.bias = true;
                break;
        case TGSI_OPCODE_TXL:
        case TGSI_OPCODE_TXL2:
-               infix = ".l";
+               args.lod = true;
                break;
        case TGSI_OPCODE_TXD:
-               infix = ".d";
+               args.deriv = true;
                break;
        case TGSI_OPCODE_TG4:
-               name = "llvm.SI.gather4";
-               infix = ".lz";
+               args.opcode = ac_image_gather4;
+               args.level_zero = true;
                break;
        default:
                assert(0);
                return;
        }
 
-       /* Add the type and suffixes .c, .o if needed. */
-       ac_build_type_name_for_intr(LLVMTypeOf(emit_data->args[0]), type, sizeof(type));
-       sprintf(intr_name, "%s%s%s%s.%s",
-               name, is_shadow ? ".c" : "", infix,
-               has_offset ? ".o" : "", type);
-
        /* 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.
@@ -4777,21 +4933,12 @@ static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
                assert(inst->Src[src_idx].Register.File == TGSI_FILE_SAMPLER);
 
                if (info->sampler_type[sampler] == TGSI_RETURN_TYPE_SINT ||
-                   info->sampler_type[sampler] == TGSI_RETURN_TYPE_UINT) {
-                       /* Texture coordinates start after:
-                        *   {offset, bias, z-compare, derivatives}
-                        * Only the offset and z-compare can occur here.
-                        */
-                       si_lower_gather4_integer(ctx, emit_data, intr_name,
-                                                (int)has_offset + (int)is_shadow);
-                       return;
-               }
+                   info->sampler_type[sampler] == TGSI_RETURN_TYPE_UINT)
+                       si_lower_gather4_integer(ctx, &args, target);
        }
 
-       emit_data->output[emit_data->chan] = lp_build_intrinsic(
-               base->gallivm->builder, intr_name, emit_data->dst_type,
-               emit_data->args, emit_data->arg_count,
-               LP_FUNC_ATTR_READNONE | LP_FUNC_ATTR_LEGACY);
+       emit_data->output[emit_data->chan] =
+               ac_build_image_opcode(&ctx->ac, &args);
 }
 
 static void si_llvm_emit_txqs(
@@ -4800,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;
@@ -4812,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;
@@ -4829,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;
@@ -4846,8 +4993,8 @@ static void si_llvm_emit_ddxy(
        idx = (opcode == TGSI_OPCODE_DDX || opcode == TGSI_OPCODE_DDX_FINE) ? 1 : 2;
 
        val = LLVMBuildBitCast(gallivm->builder, emit_data->args[0], ctx->i32, "");
-       val = ac_emit_ddxy(&ctx->ac, ctx->screen->has_ds_bpermute,
-                          mask, idx, ctx->lds, val);
+       val = ac_build_ddxy(&ctx->ac, ctx->screen->has_ds_bpermute,
+                           mask, idx, ctx->lds, val);
        emit_data->output[emit_data->chan] = val;
 }
 
@@ -4861,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;
 
@@ -4880,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) {
@@ -4895,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.
@@ -4908,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;
        }
@@ -4925,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;
@@ -4954,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) {
@@ -4970,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,
@@ -4991,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++) {
@@ -4999,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)
 {
@@ -5044,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);
@@ -5070,17 +5354,16 @@ 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));
 
-               lp_build_intrinsic(gallivm->builder, "llvm.AMDGPU.kill",
-                                  ctx->voidt, &kill, 1, LP_FUNC_ATTR_LEGACY);
+               ac_build_kill(&ctx->ac, kill);
        } else {
                lp_build_if(&if_state, gallivm, can_emit);
        }
@@ -5096,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);
@@ -5105,24 +5388,22 @@ static void si_llvm_emit_vertex(
 
                        out_val = LLVMBuildBitCast(gallivm->builder, out_val, ctx->i32, "");
 
-                       ac_build_tbuffer_store(&ctx->ac,
-                                              ctx->gsvs_ring[stream],
-                                              out_val, 1,
-                                              voffset, soffset, 0,
-                                              V_008F0C_BUF_DATA_FORMAT_32,
-                                              V_008F0C_BUF_NUM_FORMAT_UINT,
-                                              1, 0, 1, 1, 0);
+                       ac_build_buffer_store_dword(&ctx->ac,
+                                                   ctx->gsvs_ring[stream],
+                                                   out_val, 1,
+                                                   voffset, soffset, 0,
+                                                   1, 1, true, true);
                }
        }
 
        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]);
 
        /* Signal vertex emission */
-       ac_emit_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
-                       LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));       
+       ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
+                        LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
        if (!use_kill)
                lp_build_endif(&if_state);
 }
@@ -5138,8 +5419,8 @@ static void si_llvm_emit_primitive(
 
        /* Signal primitive cut */
        stream = si_llvm_get_stream(bld_base, emit_data);
-       ac_emit_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8),
-                       LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
+       ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8),
+                        LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
 }
 
 static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
@@ -5147,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
@@ -5163,7 +5444,7 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
        lp_build_intrinsic(gallivm->builder,
                           HAVE_LLVM >= 0x0309 ? "llvm.amdgcn.s.barrier"
                                               : "llvm.AMDGPU.barrier.local",
-                          ctx->voidt, NULL, 0, 0);
+                          ctx->voidt, NULL, 0, LP_FUNC_ATTR_CONVERGENT);
 }
 
 static const struct lp_build_tgsi_action tex_action = {
@@ -5201,7 +5482,8 @@ 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_attr_dereferenceable(P, UINT64_MAX);
+                       lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_NOALIAS);
+                       ac_add_attr_dereferenceable(P, UINT64_MAX);
                } else
                        lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_INREG);
        }
@@ -5278,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");
 }
@@ -5307,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];
@@ -5572,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,
@@ -5586,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);
@@ -5634,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), "");
@@ -5675,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];
 
@@ -5688,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. */
@@ -5702,8 +5980,7 @@ static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx,
        /* The intrinsic kills the thread if arg < 0. */
        bit = LLVMBuildSelect(builder, bit, LLVMConstReal(ctx->f32, 0),
                              LLVMConstReal(ctx->f32, -1), "");
-       lp_build_intrinsic(builder, "llvm.AMDGPU.kill", ctx->voidt, &bit, 1,
-                          LP_FUNC_ATTR_LEGACY);
+       ac_build_kill(&ctx->ac, bit);
 }
 
 void si_shader_binary_read_config(struct ac_shader_binary *binary,
@@ -5847,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,
@@ -6176,7 +6459,6 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
        struct lp_build_context *uint = &bld_base->uint_bld;
        struct si_shader_output_values *outputs;
        struct tgsi_shader_info *gsinfo = &gs_selector->info;
-       LLVMValueRef args[9];
        int i, r;
 
        outputs = MALLOC(gsinfo->num_outputs * sizeof(outputs[0]));
@@ -6202,17 +6484,9 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
        create_function(&ctx);
        preload_ring_buffers(&ctx);
 
-       args[0] = ctx.gsvs_ring[0];
-       args[1] = lp_build_mul_imm(uint,
-                                  LLVMGetParam(ctx.main_fn,
-                                               ctx.param_vertex_id),
-                                  4);
-       args[3] = uint->zero;
-       args[4] = uint->one;  /* OFFEN */
-       args[5] = uint->zero; /* IDXEN */
-       args[6] = uint->one;  /* GLC */
-       args[7] = uint->one;  /* SLC */
-       args[8] = uint->zero; /* TFE */
+       LLVMValueRef voffset =
+               lp_build_mul_imm(uint, LLVMGetParam(ctx.main_fn,
+                                                   ctx.param_vertex_id), 4);
 
        /* Fetch the vertex stream ID.*/
        LLVMValueRef stream_id;
@@ -6220,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) {
@@ -6250,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 */
@@ -6263,19 +6537,15 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
                                        continue;
                                }
 
-                               args[2] = lp_build_const_int32(
-                                       gallivm,
-                                       offset * gs_selector->gs_max_out_vertices * 16 * 4);
+                               LLVMValueRef soffset = LLVMConstInt(ctx.i32,
+                                       offset * gs_selector->gs_max_out_vertices * 16 * 4, 0);
                                offset++;
 
                                outputs[i].values[chan] =
-                                       LLVMBuildBitCast(gallivm->builder,
-                                                lp_build_intrinsic(gallivm->builder,
-                                                                   "llvm.SI.buffer.load.dword.i32.i32",
-                                                                   ctx.i32, args, 9,
-                                                                   LP_FUNC_ATTR_READONLY |
-                                                                   LP_FUNC_ATTR_LEGACY),
-                                                ctx.f32, "");
+                                       ac_build_buffer_load(&ctx.ac,
+                                                            ctx.gsvs_ring[0], 1,
+                                                            ctx.i32_0, voffset,
+                                                            soffset, 0, 1, 1, true);
                        }
                }
 
@@ -6299,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) {
@@ -6423,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;
@@ -6469,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)
@@ -6490,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, 5 + i);
+               LLVMValueRef p = LLVMGetOperand(inst, EXP_OUT0 + i);
 
                /* It's a constant expression. Undef outputs are eliminated too. */
                if (LLVMIsUndef(p)) {
@@ -6574,10 +6861,12 @@ static void si_eliminate_const_vs_outputs(struct si_shader_context *ctx)
                        unsigned num_args = LLVMCountParams(callee);
 
                        /* Check if this is an export instruction. */
-                       if (num_args != 9 || strcmp(name, "llvm.SI.export"))
+                       if ((num_args != 9 && num_args != 8) ||
+                           (strcmp(name, "llvm.SI.export") &&
+                            strcmp(name, "llvm.amdgcn.exp.f32")))
                                continue;
 
-                       LLVMValueRef arg = LLVMGetOperand(cur, 3);
+                       LLVMValueRef arg = LLVMGetOperand(cur, EXP_TARGET);
                        unsigned target = LLVMConstIntGetZExtValue(arg);
 
                        if (target < V_008DFC_SQ_EXP_PARAM)
@@ -6619,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;
@@ -6709,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, "");
                }
        }
@@ -7189,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;
 
@@ -7209,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)) {
@@ -7302,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 &&
@@ -7434,12 +7721,12 @@ si_get_shader_part(struct si_screen *sscreen,
 {
        struct si_shader_part *result;
 
-       pipe_mutex_lock(sscreen->shader_parts_mutex);
+       mtx_lock(&sscreen->shader_parts_mutex);
 
        /* 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;
                }
        }
@@ -7493,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;
 }
 
@@ -7640,7 +7927,7 @@ static void si_build_vs_epilog_function(struct si_shader_context *ctx,
                args.out[2] = base->undef; /* Z */
                args.out[3] = base->undef; /* W */
 
-               ac_emit_export(&ctx->ac, &args);
+               ac_build_export(&ctx->ac, &args);
        }
 
        LLVMBuildRetVoid(gallivm->builder);