radeonsi: change the bit-packing of LS out/TCS in data
[mesa.git] / src / gallium / drivers / radeonsi / si_shader.c
index f404273243f2bf5e336c5ad38c47a4d8c80c7b36..5c17c640a3c9462500594b3f040eea3f6c376a9b 100644 (file)
 #include "gallivm/lp_bld_arit.h"
 #include "gallivm/lp_bld_flow.h"
 #include "gallivm/lp_bld_misc.h"
-#include "radeon/radeon_elf_util.h"
 #include "util/u_memory.h"
 #include "util/u_string.h"
 #include "tgsi/tgsi_build.h"
 #include "tgsi/tgsi_util.h"
 #include "tgsi/tgsi_dump.h"
 
+#include "ac_binary.h"
 #include "ac_llvm_util.h"
 #include "si_shader_internal.h"
 #include "si_pipe.h"
@@ -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,
@@ -98,14 +100,6 @@ enum {
        LOCAL_ADDR_SPACE = 3,
 };
 
-#define SENDMSG_GS 2
-#define SENDMSG_GS_DONE 3
-
-#define SENDMSG_GS_OP_NOP      (0 << 4)
-#define SENDMSG_GS_OP_CUT      (1 << 4)
-#define SENDMSG_GS_OP_EMIT     (2 << 4)
-#define SENDMSG_GS_OP_EMIT_CUT (3 << 4)
-
 /**
  * Returns a unique index for a semantic name and index. The index must be
  * less than 64, so that a 64-bit bitmask of used inputs or outputs can be
@@ -181,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;
@@ -233,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;
@@ -307,82 +301,37 @@ get_tcs_out_current_patch_data_offset(struct si_shader_context *ctx)
                            "");
 }
 
-static LLVMValueRef build_gep0(struct si_shader_context *ctx,
-                              LLVMValueRef base_ptr, LLVMValueRef index)
-{
-       LLVMValueRef indices[2] = {
-               LLVMConstInt(ctx->i32, 0, 0),
-               index,
-       };
-       return LLVMBuildGEP(ctx->gallivm.builder, base_ptr,
-                           indices, 2, "");
-}
-
-static void build_indexed_store(struct si_shader_context *ctx,
-                               LLVMValueRef base_ptr, LLVMValueRef index,
-                               LLVMValueRef value)
-{
-       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-
-       LLVMBuildStore(gallivm->builder, value,
-                      build_gep0(ctx, base_ptr, index));
-}
-
-/**
- * Build an LLVM bytecode indexed load using LLVMBuildGEP + LLVMBuildLoad.
- * It's equivalent to doing a load from &base_ptr[index].
- *
- * \param base_ptr  Where the array starts.
- * \param index     The element index into the array.
- * \param uniform   Whether the base_ptr and index can be assumed to be
- *                  dynamically uniform
- */
-static LLVMValueRef build_indexed_load(struct si_shader_context *ctx,
-                                      LLVMValueRef base_ptr, LLVMValueRef index,
-                                      bool uniform)
-{
-       struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef pointer;
-
-       pointer = build_gep0(ctx, base_ptr, index);
-       if (uniform)
-               LLVMSetMetadata(pointer, ctx->uniform_md_kind, ctx->empty_md);
-       return LLVMBuildLoad(gallivm->builder, pointer, "");
-}
-
-/**
- * Do a load from &base_ptr[index], but also add a flag that it's loading
- * a constant from a dynamically uniform index.
- */
-static LLVMValueRef build_indexed_load_const(
-       struct si_shader_context *ctx,
-       LLVMValueRef base_ptr, LLVMValueRef index)
-{
-       LLVMValueRef result = build_indexed_load(ctx, base_ptr, index, true);
-       LLVMSetMetadata(result, ctx->invariant_load_md_kind, ctx->empty_md);
-       return result;
-}
-
 static LLVMValueRef get_instance_index_for_fetch(
-       struct si_shader_context *radeon_bld,
+       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
+ * to float. */
+static LLVMValueRef extract_double_to_float(struct si_shader_context *ctx,
+                                           LLVMValueRef vec4,
+                                           unsigned double_index)
+{
+       LLVMBuilderRef builder = ctx->gallivm.builder;
+       LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->gallivm.context);
+       LLVMValueRef dvec2 = LLVMBuildBitCast(builder, vec4,
+                                             LLVMVectorType(f64, 2), "");
+       LLVMValueRef index = LLVMConstInt(ctx->i32, double_index, 0);
+       LLVMValueRef value = LLVMBuildExtractElement(builder, dvec2, index, "");
+       return LLVMBuildFPTrunc(builder, value, ctx->f32, "");
 }
 
 static void declare_input_vs(
@@ -391,50 +340,72 @@ 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;
+       unsigned num_fetches;
+       unsigned fetch_stride;
 
        LLVMValueRef t_list_ptr;
        LLVMValueRef t_offset;
        LLVMValueRef t_list;
-       LLVMValueRef attribute_offset;
-       LLVMValueRef buffer_index;
-       LLVMValueRef args[3];
-       LLVMValueRef input;
+       LLVMValueRef vertex_index;
+       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 = build_indexed_load_const(ctx, t_list_ptr, t_offset);
+       t_list = ac_build_indexed_load_const(&ctx->ac, t_list_ptr, t_offset);
 
-       /* Build the attribute offset */
-       attribute_offset = lp_build_const_int32(gallivm, 0);
-
-       buffer_index = LLVMGetParam(ctx->main_fn,
+       vertex_index = LLVMGetParam(ctx->main_fn,
                                    ctx->param_vertex_index0 +
                                    input_index);
 
-       args[0] = t_list;
-       args[1] = attribute_offset;
-       args[2] = buffer_index;
-       input = lp_build_intrinsic(gallivm->builder,
-               "llvm.SI.vs.load.input", ctx->v4f32, args, 3,
-               LP_FUNC_ATTR_READNONE);
+       fix_fetch = ctx->shader->key.mono.vs.fix_fetch[input_index];
+
+       /* Do multiple loads for special formats. */
+       switch (fix_fetch) {
+       case SI_FIX_FETCH_RGB_64_FLOAT:
+               num_fetches = 3; /* 3 2-dword loads */
+               fetch_stride = 8;
+               break;
+       case SI_FIX_FETCH_RGBA_64_FLOAT:
+               num_fetches = 2; /* 2 4-dword loads */
+               fetch_stride = 16;
+               break;
+       case SI_FIX_FETCH_RGB_8:
+       case SI_FIX_FETCH_RGB_8_INT:
+               num_fetches = 3;
+               fetch_stride = 1;
+               break;
+       case SI_FIX_FETCH_RGB_16:
+       case SI_FIX_FETCH_RGB_16_INT:
+               num_fetches = 3;
+               fetch_stride = 2;
+               break;
+       default:
+               num_fetches = 1;
+               fetch_stride = 0;
+       }
+
+       for (unsigned i = 0; i < num_fetches; i++) {
+               LLVMValueRef voffset = LLVMConstInt(ctx->i32, fetch_stride * i, 0);
+
+               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, llvm_chan, "");
+                                                   input[0], llvm_chan, "");
        }
 
-       fix_fetch = (ctx->shader->key.mono.vs.fix_fetch >> (4 * input_index)) & 0xf;
-
        switch (fix_fetch) {
        case SI_FIX_FETCH_A2_SNORM:
        case SI_FIX_FETCH_A2_SSCALED:
@@ -530,6 +501,42 @@ static void declare_input_vs(
                                                    out[chan], ctx->f32, "");
                }
                break;
+       case SI_FIX_FETCH_RG_64_FLOAT:
+               for (chan = 0; chan < 2; chan++)
+                       out[chan] = extract_double_to_float(ctx, input[0], chan);
+
+               out[2] = LLVMConstReal(ctx->f32, 0);
+               out[3] = LLVMConstReal(ctx->f32, 1);
+               break;
+       case SI_FIX_FETCH_RGB_64_FLOAT:
+               for (chan = 0; chan < 3; chan++)
+                       out[chan] = extract_double_to_float(ctx, input[chan], 0);
+
+               out[3] = LLVMConstReal(ctx->f32, 1);
+               break;
+       case SI_FIX_FETCH_RGBA_64_FLOAT:
+               for (chan = 0; chan < 4; chan++) {
+                       out[chan] = extract_double_to_float(ctx, input[chan / 2],
+                                                           chan % 2);
+               }
+               break;
+       case SI_FIX_FETCH_RGB_8:
+       case SI_FIX_FETCH_RGB_8_INT:
+       case SI_FIX_FETCH_RGB_16:
+       case SI_FIX_FETCH_RGB_16_INT:
+               for (chan = 0; chan < 3; chan++) {
+                       out[chan] = LLVMBuildExtractElement(gallivm->builder,
+                                                           input[chan],
+                                                           ctx->i32_0, "");
+               }
+               if (fix_fetch == SI_FIX_FETCH_RGB_8 ||
+                   fix_fetch == SI_FIX_FETCH_RGB_16) {
+                       out[3] = LLVMConstReal(ctx->f32, 1);
+               } else {
+                       out[3] = LLVMBuildBitCast(gallivm->builder, ctx->i32_1,
+                                                 ctx->f32, "");
+               }
+               break;
        }
 }
 
@@ -539,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:
@@ -556,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;
        }
 }
 
@@ -568,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;
 }
 
@@ -592,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);
@@ -608,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;
@@ -636,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,
@@ -671,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 {
@@ -681,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
@@ -703,10 +710,11 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
  * Note that every attribute has 4 components.
  */
 static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx,
+                                              LLVMValueRef rel_patch_id,
                                                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;
 
@@ -715,9 +723,9 @@ 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, get_rel_patch_id(ctx),
+               base_addr = LLVMBuildMul(gallivm->builder, rel_patch_id,
                                         vertices_per_patch, "");
 
                base_addr = LLVMBuildAdd(gallivm->builder, base_addr,
@@ -725,7 +733,7 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx,
 
                param_stride = total_vertices;
        } else {
-               base_addr = get_rel_patch_id(ctx);
+               base_addr = rel_patch_id;
                param_stride = num_patches;
        }
 
@@ -750,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;
@@ -766,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. */
@@ -795,196 +802,52 @@ 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, vertex_index, param_index);
-}
-
-/* TBUFFER_STORE_FORMAT_{X,XY,XYZ,XYZW} <- the suffix is selected by num_channels=1..4.
- * The type of vdata must be one of i32 (num_channels=1), v2i32 (num_channels=2),
- * or v4i32 (num_channels=3,4). */
-static void build_tbuffer_store(struct si_shader_context *ctx,
-                               LLVMValueRef rsrc,
-                               LLVMValueRef vdata,
-                               unsigned num_channels,
-                               LLVMValueRef vaddr,
-                               LLVMValueRef soffset,
-                               unsigned inst_offset,
-                               unsigned dfmt,
-                               unsigned nfmt,
-                               unsigned offen,
-                               unsigned idxen,
-                               unsigned glc,
-                               unsigned slc,
-                               unsigned tfe)
-{
-       struct gallivm_state *gallivm = &ctx->gallivm;
-       LLVMValueRef args[] = {
-               rsrc,
-               vdata,
-               LLVMConstInt(ctx->i32, num_channels, 0),
-               vaddr,
-               soffset,
-               LLVMConstInt(ctx->i32, inst_offset, 0),
-               LLVMConstInt(ctx->i32, dfmt, 0),
-               LLVMConstInt(ctx->i32, nfmt, 0),
-               LLVMConstInt(ctx->i32, offen, 0),
-               LLVMConstInt(ctx->i32, idxen, 0),
-               LLVMConstInt(ctx->i32, glc, 0),
-               LLVMConstInt(ctx->i32, slc, 0),
-               LLVMConstInt(ctx->i32, tfe, 0)
-       };
-
-       /* The instruction offset field has 12 bits */
-       assert(offen || inst_offset < (1 << 12));
-
-       /* The intrinsic is overloaded, we need to add a type suffix for overloading to work. */
-       unsigned func = CLAMP(num_channels, 1, 3) - 1;
-       const char *types[] = {"i32", "v2i32", "v4i32"};
-       char name[256];
-       snprintf(name, sizeof(name), "llvm.SI.tbuffer.store.%s", types[func]);
-
-       lp_build_intrinsic(gallivm->builder, name, ctx->voidt,
-                          args, ARRAY_SIZE(args), 0);
-}
-
-static void build_tbuffer_store_dwords(struct si_shader_context *ctx,
-                                    LLVMValueRef rsrc,
-                                    LLVMValueRef vdata,
-                                    unsigned num_channels,
-                                    LLVMValueRef vaddr,
-                                    LLVMValueRef soffset,
-                                    unsigned inst_offset)
-{
-       static unsigned dfmt[] = {
-               V_008F0C_BUF_DATA_FORMAT_32,
-               V_008F0C_BUF_DATA_FORMAT_32_32,
-               V_008F0C_BUF_DATA_FORMAT_32_32_32,
-               V_008F0C_BUF_DATA_FORMAT_32_32_32_32
-       };
-       assert(num_channels >= 1 && num_channels <= 4);
-
-       build_tbuffer_store(ctx, rsrc, vdata, num_channels, vaddr, soffset,
-                           inst_offset, dfmt[num_channels-1],
-                           V_008F0C_BUF_NUM_FORMAT_UINT, 1, 0, 1, 1, 0);
-}
-
-static LLVMValueRef build_buffer_load(struct si_shader_context *ctx,
-                                      LLVMValueRef rsrc,
-                                      int num_channels,
-                                      LLVMValueRef vindex,
-                                      LLVMValueRef voffset,
-                                      LLVMValueRef soffset,
-                                      unsigned inst_offset,
-                                      unsigned glc,
-                                      unsigned slc)
-{
-       struct gallivm_state *gallivm = &ctx->gallivm;
-       unsigned func = CLAMP(num_channels, 1, 3) - 1;
-
-       if (HAVE_LLVM >= 0x309) {
-               LLVMValueRef args[] = {
-                       LLVMBuildBitCast(gallivm->builder, rsrc, ctx->v4i32, ""),
-                       vindex ? vindex : LLVMConstInt(ctx->i32, 0, 0),
-                       LLVMConstInt(ctx->i32, inst_offset, 0),
-                       LLVMConstInt(ctx->i1, glc, 0),
-                       LLVMConstInt(ctx->i1, slc, 0)
-               };
-
-               LLVMTypeRef types[] = {ctx->f32, LLVMVectorType(ctx->f32, 2),
-                                      ctx->v4f32};
-               const char *type_names[] = {"f32", "v2f32", "v4f32"};
-               char name[256];
-
-               if (voffset) {
-                       args[2] = LLVMBuildAdd(gallivm->builder, args[2], voffset,
-                                              "");
-               }
-
-               if (soffset) {
-                       args[2] = LLVMBuildAdd(gallivm->builder, args[2], soffset,
-                                              "");
-               }
-
-               snprintf(name, sizeof(name), "llvm.amdgcn.buffer.load.%s",
-                        type_names[func]);
-
-               return lp_build_intrinsic(gallivm->builder, name, types[func], args,
-                                         ARRAY_SIZE(args), LP_FUNC_ATTR_READONLY);
-       } else {
-               LLVMValueRef args[] = {
-                       LLVMBuildBitCast(gallivm->builder, rsrc, ctx->v16i8, ""),
-                       voffset ? voffset : vindex,
-                       soffset,
-                       LLVMConstInt(ctx->i32, inst_offset, 0),
-                       LLVMConstInt(ctx->i32, voffset ? 1 : 0, 0), // offen
-                       LLVMConstInt(ctx->i32, vindex ? 1 : 0, 0), //idxen
-                       LLVMConstInt(ctx->i32, glc, 0),
-                       LLVMConstInt(ctx->i32, slc, 0),
-                       LLVMConstInt(ctx->i32, 0, 0), // TFE
-               };
-
-               LLVMTypeRef types[] = {ctx->i32, LLVMVectorType(ctx->i32, 2),
-                                      ctx->v4i32};
-               const char *type_names[] = {"i32", "v2i32", "v4i32"};
-               const char *arg_type = "i32";
-               char name[256];
-
-               if (voffset && vindex) {
-                       LLVMValueRef vaddr[] = {vindex, voffset};
-
-                       arg_type = "v2i32";
-                       args[1] = lp_build_gather_values(gallivm, vaddr, 2);
-               }
-
-               snprintf(name, sizeof(name), "llvm.SI.buffer.load.dword.%s.%s",
-                        type_names[func], arg_type);
-
-               return lp_build_intrinsic(gallivm->builder, name, types[func], args,
-                                         ARRAY_SIZE(args), LP_FUNC_ATTR_READONLY);
-       }
+       return get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx),
+                                         vertex_index, param_index);
 }
 
 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 = build_buffer_load(ctx, buffer, 4, NULL, base, offset,
-                                         0, 1, 0);
+               value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
+                                            0, 1, 0, readonly_memory);
 
                return LLVMBuildBitCast(gallivm->builder, value, vec_type, "");
        }
 
        if (!tgsi_type_is_64bit(type)) {
-               value = build_buffer_load(ctx, buffer, 4, NULL, base, offset,
-                                         0, 1, 0);
+               value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
+                                            0, 1, 0, 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 = build_buffer_load(ctx, buffer, 1, NULL, base, offset,
-                                 swizzle * 4, 1, 0);
+       value = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
+                                 swizzle * 4, 1, 0, readonly_memory);
 
-       value2 = build_buffer_load(ctx, buffer, 1, NULL, base, offset,
-                                  swizzle * 4 + 4, 1, 0);
+       value2 = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
+                                  swizzle * 4 + 4, 1, 0, readonly_memory);
 
        return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
 }
@@ -1001,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) {
@@ -1010,19 +873,19 @@ 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 = build_indexed_load(ctx, ctx->lds, dw_addr, false);
+       value = ac_build_indexed_load(&ctx->ac, ctx->lds, dw_addr, false);
        if (tgsi_type_is_64bit(type)) {
                LLVMValueRef value2;
                dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
-                                      lp_build_const_int32(gallivm, 1));
-               value2 = build_indexed_load(ctx, ctx->lds, dw_addr, false);
+                                      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);
        }
 
@@ -1042,14 +905,14 @@ 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, "");
-       build_indexed_store(ctx, ctx->lds,
-                           dw_addr, value);
+       ac_build_indexed_store(&ctx->ac, ctx->lds,
+                              dw_addr, value);
 }
 
 static LLVMValueRef fetch_input_tcs(
@@ -1060,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);
 
@@ -1093,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 = build_indexed_load_const(ctx, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
+       buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
+                       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,
@@ -1113,12 +975,15 @@ 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;
        LLVMValueRef dw_addr, stride;
        LLVMValueRef rw_buffers, buffer, base, buf_addr;
        LLVMValueRef values[4];
+       bool skip_lds_store;
+       bool is_tess_factor = false;
 
        /* Only handle per-patch and per-vertex outputs here.
         * Vectors will be lowered to scalars and this function will be called again.
@@ -1133,15 +998,28 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
                stride = unpack_param(ctx, SI_PARAM_TCS_OUT_LAYOUT, 13, 8);
                dw_addr = get_tcs_out_current_patch_offset(ctx);
                dw_addr = get_dw_address(ctx, reg, NULL, stride, dw_addr);
+               skip_lds_store = !sh_info->reads_pervertex_outputs;
        } else {
                dw_addr = get_tcs_out_current_patch_data_offset(ctx);
                dw_addr = get_dw_address(ctx, reg, NULL, NULL, dw_addr);
+               skip_lds_store = !sh_info->reads_perpatch_outputs;
+
+               if (!reg->Register.Indirect) {
+                       int name = sh_info->output_semantic_name[reg->Register.Index];
+
+                       /* Always write tess factors into LDS for the TCS epilog. */
+                       if (name == TGSI_SEMANTIC_TESSINNER ||
+                           name == TGSI_SEMANTIC_TESSOUTER) {
+                               skip_lds_store = false;
+                               is_tess_factor = true;
+                       }
+               }
        }
 
        rw_buffers = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_RW_BUFFERS);
-       buffer = build_indexed_load_const(ctx, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
+       buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
+                       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);
@@ -1151,25 +1029,27 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
                LLVMValueRef value = dst[chan_index];
 
                if (inst->Instruction.Saturate)
-                       value = si_llvm_saturate(bld_base, value);
+                       value = ac_build_clamp(&ctx->ac, value);
 
-               lds_store(bld_base, chan_index, dw_addr, value);
+               /* Skip LDS stores if there is no LDS read of this output. */
+               if (!skip_lds_store)
+                       lds_store(bld_base, chan_index, dw_addr, value);
 
                value = LLVMBuildBitCast(gallivm->builder, value, ctx->i32, "");
                values[chan_index] = value;
 
-               if (inst->Dst[0].Register.WriteMask != 0xF) {
-                       build_tbuffer_store_dwords(ctx, buffer, value, 1,
-                                                  buf_addr, base,
-                                                  4 * chan_index);
+               if (inst->Dst[0].Register.WriteMask != 0xF && !is_tess_factor) {
+                       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) {
-               LLVMValueRef value = lp_build_gather_values(bld_base->base.gallivm,
+       if (inst->Dst[0].Register.WriteMask == 0xF && !is_tess_factor) {
+               LLVMValueRef value = lp_build_gather_values(gallivm,
                                                            values, 4);
-               build_tbuffer_store_dwords(ctx, 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);
        }
 }
 
@@ -1179,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];
@@ -1205,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);
        }
 
@@ -1223,27 +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);
+       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);
+               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);
        }
@@ -1281,80 +1149,6 @@ static int lookup_interp_param_index(unsigned interpolate, unsigned location)
        }
 }
 
-static LLVMValueRef build_fs_interp(
-       struct lp_build_tgsi_context *bld_base,
-       LLVMValueRef llvm_chan,
-       LLVMValueRef attr_number,
-       LLVMValueRef params,
-       LLVMValueRef i,
-       LLVMValueRef j) {
-
-       struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef args[5];
-       LLVMValueRef p1;
-       if (HAVE_LLVM < 0x0400) {
-               LLVMValueRef ij[2];
-               ij[0] = LLVMBuildBitCast(gallivm->builder, i, ctx->i32, "");
-               ij[1] = LLVMBuildBitCast(gallivm->builder, j, ctx->i32, "");
-
-               args[0] = llvm_chan;
-               args[1] = attr_number;
-               args[2] = params;
-               args[3] = lp_build_gather_values(gallivm, ij, 2);
-               return lp_build_intrinsic(gallivm->builder, "llvm.SI.fs.interp",
-                                         ctx->f32, args, 4,
-                                         LP_FUNC_ATTR_READNONE);
-       }
-
-       args[0] = i;
-       args[1] = llvm_chan;
-       args[2] = attr_number;
-       args[3] = params;
-
-       p1 = lp_build_intrinsic(gallivm->builder, "llvm.amdgcn.interp.p1",
-                               ctx->f32, args, 4, LP_FUNC_ATTR_READNONE);
-
-       args[0] = p1;
-       args[1] = j;
-       args[2] = llvm_chan;
-       args[3] = attr_number;
-       args[4] = params;
-
-       return lp_build_intrinsic(gallivm->builder, "llvm.amdgcn.interp.p2",
-                                 ctx->f32, args, 5, LP_FUNC_ATTR_READNONE);
-}
-
-static LLVMValueRef build_fs_interp_mov(
-       struct lp_build_tgsi_context *bld_base,
-       LLVMValueRef parameter,
-       LLVMValueRef llvm_chan,
-       LLVMValueRef attr_number,
-       LLVMValueRef params) {
-
-       struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef args[4];
-       if (HAVE_LLVM < 0x0400) {
-               args[0] = llvm_chan;
-               args[1] = attr_number;
-               args[2] = params;
-
-               return lp_build_intrinsic(gallivm->builder,
-                                         "llvm.SI.fs.constant",
-                                         ctx->f32, args, 3,
-                                         LP_FUNC_ATTR_READNONE);
-       }
-
-       args[0] = parameter;
-       args[1] = llvm_chan;
-       args[2] = attr_number;
-       args[3] = params;
-
-       return lp_build_intrinsic(gallivm->builder, "llvm.amdgcn.interp.mov",
-                                 ctx->f32, args, 4, LP_FUNC_ATTR_READNONE);
-}
-
 /**
  * Interpolate a fragment shader input.
  *
@@ -1380,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;
 
@@ -1404,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 &&
@@ -1428,28 +1219,28 @@ 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) {
-                               front = build_fs_interp(bld_base, llvm_chan,
+                               front = ac_build_fs_interp(&ctx->ac, llvm_chan,
                                                        attr_number, prim_mask,
                                                        i, j);
-                               back = build_fs_interp(bld_base, llvm_chan,
+                               back = ac_build_fs_interp(&ctx->ac, llvm_chan,
                                                        back_attr_number, prim_mask,
                                                        i, j);
                        } else {
-                               front = build_fs_interp_mov(bld_base,
-                                       lp_build_const_int32(gallivm, 2), /* P0 */
+                               front = ac_build_fs_interp_mov(&ctx->ac,
+                                       LLVMConstInt(ctx->i32, 2, 0), /* P0 */
                                        llvm_chan, attr_number, prim_mask);
-                               back = build_fs_interp_mov(bld_base,
-                                       lp_build_const_int32(gallivm, 2), /* P0 */
+                               back = ac_build_fs_interp_mov(&ctx->ac,
+                                       LLVMConstInt(ctx->i32, 2, 0), /* P0 */
                                        llvm_chan, back_attr_number, prim_mask);
                        }
 
@@ -1461,26 +1252,26 @@ static void interp_fs_input(struct si_shader_context *ctx,
                }
        } else if (semantic_name == TGSI_SEMANTIC_FOG) {
                if (interp) {
-                       result[0] = build_fs_interp(bld_base, uint->zero,
-                                               attr_number, prim_mask, i, j);
+                       result[0] = ac_build_fs_interp(&ctx->ac, ctx->i32_0,
+                                                      attr_number, prim_mask, i, j);
                } else {
-                       result[0] = build_fs_interp_mov(bld_base, uint->zero,
-                               lp_build_const_int32(gallivm, 2), /* P0 */
-                               attr_number, prim_mask);
+                       result[0] = ac_build_fs_interp_mov(&ctx->ac, 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] = build_fs_interp(bld_base,
+                               result[chan] = ac_build_fs_interp(&ctx->ac,
                                        llvm_chan, attr_number, prim_mask, i, j);
                        } else {
-                               result[chan] = build_fs_interp_mov(bld_base,
-                                       lp_build_const_int32(gallivm, 2), /* P0 */
+                               result[chan] = ac_build_fs_interp_mov(&ctx->ac,
+                                       LLVMConstInt(ctx->i32, 2, 0), /* P0 */
                                        llvm_chan, attr_number, prim_mask);
                        }
                }
@@ -1488,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;
 
@@ -1537,55 +1326,11 @@ 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);
 }
 
-/**
- * Set range metadata on an instruction.  This can only be used on load and
- * call instructions.  If you know an instruction can only produce the values
- * 0, 1, 2, you would do set_range_metadata(value, 0, 3);
- * \p lo is the minimum value inclusive.
- * \p hi is the maximum value exclusive.
- */
-static void set_range_metadata(struct si_shader_context *ctx,
-                              LLVMValueRef value, unsigned lo, unsigned hi)
-{
-       LLVMValueRef range_md, md_args[2];
-       LLVMTypeRef type = LLVMTypeOf(value);
-       LLVMContextRef context = LLVMGetTypeContext(type);
-
-       md_args[0] = LLVMConstInt(type, lo, false);
-       md_args[1] = LLVMConstInt(type, hi, false);
-       range_md = LLVMMDNodeInContext(context, md_args, 2);
-       LLVMSetMetadata(value, ctx->range_md_kind, range_md);
-}
-
-static LLVMValueRef get_thread_id(struct si_shader_context *ctx)
-{
-       struct gallivm_state *gallivm = &ctx->gallivm;
-       LLVMValueRef tid;
-
-       if (HAVE_LLVM < 0x0308) {
-               tid = lp_build_intrinsic(gallivm->builder, "llvm.SI.tid",
-                               ctx->i32,   NULL, 0, LP_FUNC_ATTR_READNONE);
-       } else {
-               LLVMValueRef tid_args[2];
-               tid_args[0] = lp_build_const_int32(gallivm, 0xffffffff);
-               tid_args[1] = lp_build_const_int32(gallivm, 0);
-               tid_args[1] = lp_build_intrinsic(gallivm->builder,
-                                       "llvm.amdgcn.mbcnt.lo", ctx->i32,
-                                       tid_args, 2, LP_FUNC_ATTR_READNONE);
-
-               tid = lp_build_intrinsic(gallivm->builder,
-                                       "llvm.amdgcn.mbcnt.hi", ctx->i32,
-                                       tid_args, 2, LP_FUNC_ATTR_READNONE);
-       }
-       set_range_metadata(ctx, tid, 0, 64);
-       return tid;
-}
 
 /**
  * Load a dword from a constant buffer.
@@ -1598,76 +1343,75 @@ static LLVMValueRef buffer_load_const(struct si_shader_context *ctx,
        LLVMValueRef args[2] = {resource, offset};
 
        return lp_build_intrinsic(builder, "llvm.SI.load.const", ctx->f32, args, 2,
-                              LP_FUNC_ATTR_READNONE);
+                                 LP_FUNC_ATTR_READNONE |
+                                 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 resource = build_indexed_load_const(ctx, desc, buf_index);
+       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;
 
@@ -1675,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");
@@ -1684,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);
@@ -1696,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;
@@ -1722,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
                };
@@ -1761,15 +1505,15 @@ static void declare_system_value(
 
                rw_buffers = LLVMGetParam(ctx->main_fn,
                                        SI_PARAM_RW_BUFFERS);
-               buffer = build_indexed_load_const(ctx, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
+               buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
+                       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, NULL,
-                                         lp_build_const_int32(gallivm, param));
+               addr = get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx), NULL,
+                                         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;
        }
@@ -1780,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 = build_indexed_load_const(ctx, buf, slot);
+               buf = ac_build_indexed_load_const(&ctx->ac, buf, slot);
                offset = decl->Semantic.Name == TGSI_SEMANTIC_DEFAULT_TESSINNER_SI ? 4 : 0;
 
                for (i = 0; i < 4; i++)
                        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:
@@ -1814,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:
@@ -1845,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;
@@ -1882,7 +1664,7 @@ static LLVMValueRef load_const_buffer_desc(struct si_shader_context *ctx, int i)
        LLVMValueRef list_ptr = LLVMGetParam(ctx->main_fn,
                                             SI_PARAM_CONST_BUFFERS);
 
-       return build_indexed_load_const(ctx, list_ptr,
+       return ac_build_indexed_load_const(&ctx->ac, list_ptr,
                                        LLVMConstInt(ctx->i32, i, 0));
 }
 
@@ -1906,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;
@@ -1918,7 +1700,7 @@ static LLVMValueRef fetch_constant(
                index = get_bounded_indirect_index(ctx, &reg->DimIndirect,
                                                   reg->Dimension.Index,
                                                   SI_NUM_CONST_BUFFERS);
-               bufp = build_indexed_load_const(ctx, ptr, index);
+               bufp = ac_build_indexed_load_const(&ctx->ac, ptr, index);
        } else
                bufp = load_const_buffer_desc(ctx, buf);
 
@@ -1927,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);
        }
@@ -1950,54 +1732,52 @@ 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 */
 static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
                                     LLVMValueRef *values,
                                     unsigned target,
-                                    LLVMValueRef *args)
+                                    struct ac_export_args *args)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct lp_build_context *uint = &ctx->bld_base.uint_bld;
        struct lp_build_context *base = &bld_base->base;
-       struct gallivm_state *gallivm = base->gallivm;
-       LLVMBuilderRef builder = base->gallivm->builder;
+       LLVMBuilderRef builder = ctx->gallivm.builder;
        LLVMValueRef val[4];
        unsigned spi_shader_col_format = V_028714_SPI_SHADER_32_ABGR;
        unsigned chan;
-       bool is_int8;
+       bool is_int8, is_int10;
 
        /* Default is 0xf. Adjusted below depending on the format. */
-       args[0] = lp_build_const_int32(base->gallivm, 0xf); /* writemask */
+       args->enabled_channels = 0xf; /* writemask */
 
        /* Specify whether the EXEC mask represents the valid mask */
-       args[1] = uint->zero;
+       args->valid_mask = 0;
 
        /* Specify whether this is the last export */
-       args[2] = uint->zero;
+       args->done = 0;
 
        /* Specify the target we are exporting */
-       args[3] = lp_build_const_int32(base->gallivm, target);
+       args->target = target;
 
        if (ctx->type == PIPE_SHADER_FRAGMENT) {
                const struct si_shader_key *key = &ctx->shader->key;
@@ -2007,39 +1787,40 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
                assert(cbuf >= 0 && cbuf < 8);
                spi_shader_col_format = (col_formats >> (cbuf * 4)) & 0xf;
                is_int8 = (key->part.ps.epilog.color_is_int8 >> cbuf) & 0x1;
+               is_int10 = (key->part.ps.epilog.color_is_int10 >> cbuf) & 0x1;
        }
 
-       args[4] = uint->zero; /* COMPR flag */
-       args[5] = base->undef;
-       args[6] = base->undef;
-       args[7] = base->undef;
-       args[8] = base->undef;
+       args->compr = false;
+       args->out[0] = base->undef;
+       args->out[1] = base->undef;
+       args->out[2] = base->undef;
+       args->out[3] = base->undef;
 
        switch (spi_shader_col_format) {
        case V_028714_SPI_SHADER_ZERO:
-               args[0] = uint->zero; /* writemask */
-               args[3] = lp_build_const_int32(base->gallivm, V_008DFC_SQ_EXP_NULL);
+               args->enabled_channels = 0; /* writemask */
+               args->target = V_008DFC_SQ_EXP_NULL;
                break;
 
        case V_028714_SPI_SHADER_32_R:
-               args[0] = uint->one; /* writemask */
-               args[5] = values[0];
+               args->enabled_channels = 1; /* writemask */
+               args->out[0] = values[0];
                break;
 
        case V_028714_SPI_SHADER_32_GR:
-               args[0] = lp_build_const_int32(base->gallivm, 0x3); /* writemask */
-               args[5] = values[0];
-               args[6] = values[1];
+               args->enabled_channels = 0x3; /* writemask */
+               args->out[0] = values[0];
+               args->out[1] = values[1];
                break;
 
        case V_028714_SPI_SHADER_32_AR:
-               args[0] = lp_build_const_int32(base->gallivm, 0x9); /* writemask */
-               args[5] = values[0];
-               args[8] = values[3];
+               args->enabled_channels = 0x9; /* writemask */
+               args->out[0] = values[0];
+               args->out[3] = values[3];
                break;
 
        case V_028714_SPI_SHADER_FP16_ABGR:
-               args[4] = uint->one; /* COMPR flag */
+               args->compr = 1; /* COMPR flag */
 
                for (chan = 0; chan < 2; chan++) {
                        LLVMValueRef pack_args[2] = {
@@ -2048,32 +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);
-                       args[chan + 5] =
-                               LLVMBuildBitCast(base->gallivm->builder,
+                       packed = ac_build_cvt_pkrtz_f16(&ctx->ac, pack_args);
+                       args->out[chan] =
+                               LLVMBuildBitCast(ctx->gallivm.builder,
                                                 packed, ctx->f32, "");
                }
                break;
 
        case V_028714_SPI_SHADER_UNORM16_ABGR:
                for (chan = 0; chan < 4; chan++) {
-                       val[chan] = si_llvm_saturate(bld_base, 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[4] = uint->one; /* COMPR flag */
-               args[5] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int16(gallivm, val));
-               args[6] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int16(gallivm, val+2));
+               args->compr = 1; /* COMPR flag */
+               args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
+                                 si_llvm_pack_two_int16(ctx, val));
+               args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT,
+                                 si_llvm_pack_two_int16(ctx, val+2));
                break;
 
        case V_028714_SPI_SHADER_SNORM16_ABGR:
@@ -2081,74 +1859,83 @@ 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[4] = uint->one; /* COMPR flag */
-               args[5] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int32_as_int16(gallivm, val));
-               args[6] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int32_as_int16(gallivm, val+2));
+               args->compr = 1; /* COMPR flag */
+               args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
+                                 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(ctx, val+2));
                break;
 
        case V_028714_SPI_SHADER_UINT16_ABGR: {
-               LLVMValueRef max = lp_build_const_int32(gallivm, is_int8 ?
-                                                       255 : 65535);
+               LLVMValueRef max_rgb = LLVMConstInt(ctx->i32,
+                       is_int8 ? 255 : is_int10 ? 1023 : 65535, 0);
+               LLVMValueRef max_alpha =
+                       !is_int10 ? max_rgb : LLVMConstInt(ctx->i32, 3, 0);
+
                /* Clamp. */
                for (chan = 0; chan < 4; chan++) {
                        val[chan] = bitcast(bld_base, TGSI_TYPE_UNSIGNED, values[chan]);
                        val[chan] = lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_UMIN,
-                                                             val[chan], max);
+                                       val[chan],
+                                       chan == 3 ? max_alpha : max_rgb);
                }
 
-               args[4] = uint->one; /* COMPR flag */
-               args[5] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int16(gallivm, val));
-               args[6] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int16(gallivm, val+2));
+               args->compr = 1; /* COMPR flag */
+               args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
+                                 si_llvm_pack_two_int16(ctx, val));
+               args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT,
+                                 si_llvm_pack_two_int16(ctx, val+2));
                break;
        }
 
        case V_028714_SPI_SHADER_SINT16_ABGR: {
-               LLVMValueRef max = lp_build_const_int32(gallivm, is_int8 ?
-                                                       127 : 32767);
-               LLVMValueRef min = lp_build_const_int32(gallivm, is_int8 ?
-                                                       -128 : -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 : ctx->i32_1;
+               LLVMValueRef min_alpha =
+                       !is_int10 ? min_rgb : LLVMConstInt(ctx->i32, -2, 0);
+
                /* Clamp. */
                for (chan = 0; chan < 4; chan++) {
                        val[chan] = bitcast(bld_base, TGSI_TYPE_UNSIGNED, values[chan]);
                        val[chan] = lp_build_emit_llvm_binary(bld_base,
-                                                             TGSI_OPCODE_IMIN,
-                                                             val[chan], max);
+                                       TGSI_OPCODE_IMIN,
+                                       val[chan], chan == 3 ? max_alpha : max_rgb);
                        val[chan] = lp_build_emit_llvm_binary(bld_base,
-                                                             TGSI_OPCODE_IMAX,
-                                                             val[chan], min);
+                                       TGSI_OPCODE_IMAX,
+                                       val[chan], chan == 3 ? min_alpha : min_rgb);
                }
 
-               args[4] = uint->one; /* COMPR flag */
-               args[5] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int32_as_int16(gallivm, val));
-               args[6] = bitcast(bld_base, TGSI_TYPE_FLOAT,
-                                 si_llvm_pack_two_int32_as_int16(gallivm, val+2));
+               args->compr = 1; /* COMPR flag */
+               args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
+                                 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(ctx, val+2));
                break;
        }
 
        case V_028714_SPI_SHADER_32_ABGR:
-               memcpy(&args[5], values, sizeof(values[0]) * 4);
+               memcpy(&args->out[0], values, sizeof(values[0]) * 4);
                break;
        }
 }
@@ -2157,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,
@@ -2170,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, 0);
+               ac_build_kill(&ctx->ac, arg);
        } else {
-               lp_build_intrinsic(gallivm->builder, "llvm.AMDGPU.kilp",
-                                  ctx->voidt, NULL, 0, 0);
+               ac_build_kill(&ctx->ac, NULL);
        }
 }
 
@@ -2186,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 */
@@ -2202,56 +1986,54 @@ 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, "");
 }
 
 static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base,
-                                   LLVMValueRef (*pos)[9], LLVMValueRef *out_elts)
+                                   struct ac_export_args *pos, LLVMValueRef *out_elts)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct lp_build_context *base = &bld_base->base;
-       struct lp_build_context *uint = &ctx->bld_base.uint_bld;
        unsigned reg_index;
        unsigned chan;
        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 const_resource = build_indexed_load_const(ctx, ptr, constbuf_index);
+       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 ++) {
-               LLVMValueRef *args = pos[2 + reg_index];
+               struct ac_export_args *args = &pos[2 + reg_index];
 
-               args[5] =
-               args[6] =
-               args[7] =
-               args[8] = lp_build_const_float(base->gallivm, 0.0f);
+               args->out[0] =
+               args->out[1] =
+               args->out[2] =
+               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++) {
                        for (const_chan = 0; const_chan < TGSI_NUM_CHANNELS; const_chan++) {
-                               args[1] = lp_build_const_int32(base->gallivm,
-                                                              ((reg_index * 4 + chan) * 4 +
-                                                               const_chan) * 4);
+                               LLVMValueRef addr =
+                                       LLVMConstInt(ctx->i32, ((reg_index * 4 + chan) * 4 +
+                                                               const_chan) * 4, 0);
                                base_elt = buffer_load_const(ctx, const_resource,
-                                                            args[1]);
-                               args[5 + chan] =
-                                       lp_build_add(base, args[5 + chan],
+                                                            addr);
+                               args->out[chan] =
+                                       lp_build_add(base, args->out[chan],
                                                     lp_build_mul(base, base_elt,
                                                                  out_elts[const_chan]));
                        }
                }
 
-               args[0] = lp_build_const_int32(base->gallivm, 0xf);
-               args[1] = uint->zero;
-               args[2] = uint->zero;
-               args[3] = lp_build_const_int32(base->gallivm,
-                                              V_008DFC_SQ_EXP_POS + 2 + reg_index);
-               args[4] = uint->zero;
+               args->enabled_channels = 0xf;
+               args->valid_mask = 0;
+               args->done = 0;
+               args->target = V_008DFC_SQ_EXP_POS + 2 + reg_index;
+               args->compr = 0;
        }
 }
 
@@ -2320,11 +2102,11 @@ static void emit_streamout_output(struct si_shader_context *ctx,
                break;
        }
 
-       build_tbuffer_store_dwords(ctx, so_buffers[buf_idx],
-                                  vdata, num_comps,
-                                  so_write_offsets[buf_idx],
-                                  LLVMConstInt(ctx->i32, 0, 0),
-                                  stream_out->dst_offset * 4);
+       ac_build_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);
 }
 
 /**
@@ -2346,7 +2128,7 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
        LLVMValueRef so_vtx_count =
                unpack_param(ctx, ctx->param_streamout_config, 16, 7);
 
-       LLVMValueRef tid = get_thread_id(ctx);
+       LLVMValueRef tid = ac_get_thread_id(&ctx->ac);
 
        /* can_emit = tid < so_vtx_count; */
        LLVMValueRef can_emit =
@@ -2381,10 +2163,10 @@ 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] = build_indexed_load_const(ctx, buf_ptr, offset);
+                       so_buffers[i] = ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
 
                        LLVMValueRef so_offset = LLVMGetParam(ctx->main_fn,
                                                              ctx->param_streamout_offset[i]);
@@ -2421,9 +2203,7 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base,
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct si_shader *shader = ctx->shader;
        struct lp_build_context *base = &bld_base->base;
-       struct lp_build_context *uint = &ctx->bld_base.uint_bld;
-       LLVMValueRef args[9];
-       LLVMValueRef pos_args[4][9] = { { 0 } };
+       struct ac_export_args args, pos_args[4] = {};
        LLVMValueRef psize_value = NULL, edgeflag_value = NULL, layer_value = NULL, viewport_index_value = NULL;
        unsigned semantic_name, semantic_index;
        unsigned target;
@@ -2513,16 +2293,14 @@ handle_semantic:
                                semantic_name);
                }
 
-               si_llvm_init_export_args(bld_base, outputs[i].values, target, args);
+               si_llvm_init_export_args(bld_base, outputs[i].values, target, &args);
 
                if (target >= V_008DFC_SQ_EXP_POS &&
                    target <= (V_008DFC_SQ_EXP_POS + 3)) {
-                       memcpy(pos_args[target - V_008DFC_SQ_EXP_POS],
-                              args, sizeof(args));
+                       memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS],
+                              &args, sizeof(args));
                } else {
-                       lp_build_intrinsic(base->gallivm->builder,
-                                          "llvm.SI.export", ctx->voidt,
-                                          args, 9, 0);
+                       ac_build_export(&ctx->ac, &args);
                }
 
                if (semantic_name == TGSI_SEMANTIC_CLIPDIST) {
@@ -2534,16 +2312,16 @@ handle_semantic:
        shader->info.nr_param_exports = param_count;
 
        /* We need to add the position output manually if it's missing. */
-       if (!pos_args[0][0]) {
-               pos_args[0][0] = lp_build_const_int32(base->gallivm, 0xf); /* writemask */
-               pos_args[0][1] = uint->zero; /* EXEC mask */
-               pos_args[0][2] = uint->zero; /* last export? */
-               pos_args[0][3] = lp_build_const_int32(base->gallivm, V_008DFC_SQ_EXP_POS);
-               pos_args[0][4] = uint->zero; /* COMPR flag */
-               pos_args[0][5] = base->zero; /* X */
-               pos_args[0][6] = base->zero; /* Y */
-               pos_args[0][7] = base->zero; /* Z */
-               pos_args[0][8] = base->one;  /* W */
+       if (!pos_args[0].out[0]) {
+               pos_args[0].enabled_channels = 0xf; /* writemask */
+               pos_args[0].valid_mask = 0; /* EXEC mask */
+               pos_args[0].done = 0; /* last export? */
+               pos_args[0].target = V_008DFC_SQ_EXP_POS;
+               pos_args[0].compr = 0; /* COMPR flag */
+               pos_args[0].out[0] = base->zero; /* X */
+               pos_args[0].out[1] = base->zero; /* Y */
+               pos_args[0].out[2] = base->zero; /* Z */
+               pos_args[0].out[3] = base->one;  /* W */
        }
 
        /* Write the misc vector (point size, edgeflag, layer, viewport). */
@@ -2551,64 +2329,62 @@ handle_semantic:
            shader->selector->info.writes_edgeflag ||
            shader->selector->info.writes_viewport_index ||
            shader->selector->info.writes_layer) {
-               pos_args[1][0] = lp_build_const_int32(base->gallivm, /* writemask */
-                                                     shader->selector->info.writes_psize |
-                                                     (shader->selector->info.writes_edgeflag << 1) |
-                                                     (shader->selector->info.writes_layer << 2) |
-                                                     (shader->selector->info.writes_viewport_index << 3));
-               pos_args[1][1] = uint->zero; /* EXEC mask */
-               pos_args[1][2] = uint->zero; /* last export? */
-               pos_args[1][3] = lp_build_const_int32(base->gallivm, V_008DFC_SQ_EXP_POS + 1);
-               pos_args[1][4] = uint->zero; /* COMPR flag */
-               pos_args[1][5] = base->zero; /* X */
-               pos_args[1][6] = base->zero; /* Y */
-               pos_args[1][7] = base->zero; /* Z */
-               pos_args[1][8] = base->zero; /* W */
+               pos_args[1].enabled_channels = shader->selector->info.writes_psize |
+                                              (shader->selector->info.writes_edgeflag << 1) |
+                                              (shader->selector->info.writes_layer << 2) |
+                                              (shader->selector->info.writes_viewport_index << 3);
+               pos_args[1].valid_mask = 0; /* EXEC mask */
+               pos_args[1].done = 0; /* last export? */
+               pos_args[1].target = V_008DFC_SQ_EXP_POS + 1;
+               pos_args[1].compr = 0; /* COMPR flag */
+               pos_args[1].out[0] = base->zero; /* X */
+               pos_args[1].out[1] = base->zero; /* Y */
+               pos_args[1].out[2] = base->zero; /* Z */
+               pos_args[1].out[3] = base->zero; /* W */
 
                if (shader->selector->info.writes_psize)
-                       pos_args[1][5] = psize_value;
+                       pos_args[1].out[0] = psize_value;
 
                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][6] = LLVMBuildBitCast(base->gallivm->builder,
+                       pos_args[1].out[1] = LLVMBuildBitCast(ctx->gallivm.builder,
                                                          edgeflag_value,
                                                          ctx->f32, "");
                }
 
                if (shader->selector->info.writes_layer)
-                       pos_args[1][7] = layer_value;
+                       pos_args[1].out[2] = layer_value;
 
                if (shader->selector->info.writes_viewport_index)
-                       pos_args[1][8] = viewport_index_value;
+                       pos_args[1].out[3] = viewport_index_value;
        }
 
        for (i = 0; i < 4; i++)
-               if (pos_args[i][0])
+               if (pos_args[i].out[0])
                        shader->info.nr_pos_exports++;
 
        pos_idx = 0;
        for (i = 0; i < 4; i++) {
-               if (!pos_args[i][0])
+               if (!pos_args[i].out[0])
                        continue;
 
                /* Specify the target we are exporting */
-               pos_args[i][3] = lp_build_const_int32(base->gallivm, V_008DFC_SQ_EXP_POS + pos_idx++);
+               pos_args[i].target = V_008DFC_SQ_EXP_POS + pos_idx++;
 
                if (pos_idx == shader->info.nr_pos_exports)
                        /* Specify that this is the last export */
-                       pos_args[i][2] = uint->one;
+                       pos_args[i].done = 1;
 
-               lp_build_intrinsic(base->gallivm->builder, "llvm.SI.export",
-                                  ctx->voidt, pos_args[i], 9, 0);
+               ac_build_export(&ctx->ac, &pos_args[i]);
        }
 }
 
@@ -2619,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;
@@ -2627,12 +2403,12 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
        invocation_id = unpack_param(ctx, SI_PARAM_REL_IDS, 8, 5);
 
        rw_buffers = LLVMGetParam(ctx->main_fn, SI_PARAM_RW_BUFFERS);
-       buffer = build_indexed_load_const(ctx, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_OFFCHIP));
+       buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
+                       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);
@@ -2643,18 +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);
 
-               build_tbuffer_store_dwords(ctx, 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);
        }
 }
 
@@ -2664,11 +2441,11 @@ 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;
-       LLVMValueRef out[6], vec0, vec1, rw_buffers, tf_base;
+       LLVMValueRef out[6], vec0, vec1, rw_buffers, tf_base, inner[4], outer[4];
        unsigned stride, outer_comps, inner_comps, i;
        struct lp_build_if_state if_ctx, inner_if_ctx;
 
@@ -2682,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) {
@@ -2714,23 +2491,32 @@ 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);
+               outer[i] = LLVMGetUndef(ctx->i32);
+       }
 
        if (shader->key.part.tcs.epilog.prim_mode == PIPE_PRIM_LINES) {
                /* For isolines, the hardware expects tess factors in the
                 * reverse order from what GLSL / TGSI specify.
                 */
-               out[0] = lds_load(bld_base, TGSI_TYPE_SIGNED, 1, lds_outer);
-               out[1] = lds_load(bld_base, TGSI_TYPE_SIGNED, 0, lds_outer);
+               outer[0] = out[1] = lds_load(bld_base, TGSI_TYPE_SIGNED, 0, lds_outer);
+               outer[1] = out[0] = lds_load(bld_base, TGSI_TYPE_SIGNED, 1, lds_outer);
        } else {
-               for (i = 0; i < outer_comps; i++)
-                       out[i] = lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_outer);
-               for (i = 0; i < inner_comps; i++)
-                       out[outer_comps+i] = lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_inner);
+               for (i = 0; i < outer_comps; i++) {
+                       outer[i] = out[i] =
+                               lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_outer);
+               }
+               for (i = 0; i < inner_comps; i++) {
+                       inner[i] = out[outer_comps+i] =
+                               lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_inner);
+               }
        }
 
        /* Convert the outputs to vectors for stores. */
@@ -2743,32 +2529,71 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
        /* Get the buffer. */
        rw_buffers = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_RW_BUFFERS);
-       buffer = build_indexed_load_const(ctx, rw_buffers,
-                       lp_build_const_int32(gallivm, SI_HS_RING_TESS_FACTOR));
+       buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
+                       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. */
-       build_tbuffer_store_dwords(ctx, 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. */
-       build_tbuffer_store_dwords(ctx, 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)
-               build_tbuffer_store_dwords(ctx, 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) {
+               LLVMValueRef buf, base, inner_vec, outer_vec, tf_outer_offset;
+               LLVMValueRef tf_inner_offset;
+               unsigned param_outer, param_inner;
+
+               buf = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
+                               LLVMConstInt(ctx->i32, SI_HS_RING_TESS_OFFCHIP, 0));
+               base = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
+
+               param_outer = si_shader_io_get_unique_index(
+                                     TGSI_SEMANTIC_TESSOUTER, 0);
+               tf_outer_offset = get_tcs_tes_buffer_address(ctx, rel_patch_id, NULL,
+                                       LLVMConstInt(ctx->i32, param_outer, 0));
+
+               outer_vec = lp_build_gather_values(gallivm, outer,
+                                                  util_next_power_of_two(outer_comps));
+
+               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);
+                       tf_inner_offset = get_tcs_tes_buffer_address(ctx, rel_patch_id, NULL,
+                                       LLVMConstInt(ctx->i32, param_inner, 0));
+
+                       inner_vec = inner_comps == 1 ? inner[0] :
+                                   lp_build_gather_values(gallivm, inner, inner_comps);
+                       ac_build_buffer_store_dword(&ctx->ac, buf, inner_vec,
+                                                   inner_comps, tf_inner_offset,
+                                                   base, 0, 1, 0, true, false);
+               }
+       }
+
        lp_build_endif(&if_ctx);
 }
 
@@ -2777,6 +2602,7 @@ static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
        LLVMValueRef rel_patch_id, invocation_id, tf_lds_offset;
+       LLVMValueRef offchip_soffset, offchip_layout;
 
        si_copy_tcs_inputs(bld_base);
 
@@ -2785,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;
@@ -2796,15 +2622,22 @@ 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, "");
 
-       /* Tess factor buffer soffset is after user SGPRs. */
+       /* Tess offchip and factor buffer soffset are after user SGPRs. */
+       offchip_layout = LLVMGetParam(ctx->main_fn,
+                                     SI_PARAM_TCS_OFFCHIP_LAYOUT);
+       offchip_soffset = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
        tf_soffset = LLVMGetParam(ctx->main_fn,
                                  SI_PARAM_TESS_FACTOR_OFFSET);
+       ret = LLVMBuildInsertValue(builder, ret, offchip_layout,
+                                  SI_SGPR_TCS_OFFCHIP_LAYOUT, "");
+       ret = LLVMBuildInsertValue(builder, ret, offchip_soffset,
+                                  SI_TCS_NUM_USER_SGPR, "");
        ret = LLVMBuildInsertValue(builder, ret, tf_soffset,
                                   SI_TCS_NUM_USER_SGPR + 1, "");
 
@@ -2825,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, "");
 
@@ -2842,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,
@@ -2854,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,
@@ -2877,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, "");
 
-                       build_tbuffer_store(ctx,
-                                           ctx->esgs_ring,
-                                           out_val, 1,
-                                           LLVMGetUndef(ctx->i32), soffset,
-                                           (4 * param_index + chan) * 4,
-                                           V_008F0C_BUF_DATA_FORMAT_32,
-                                           V_008F0C_BUF_NUM_FORMAT_UINT,
-                                           0, 0, 1, 1, 0);
+                       ac_build_buffer_store_dword(&ctx->ac,
+                                                   ctx->esgs_ring,
+                                                   out_val, 1, NULL, soffset,
+                                                   (4 * param_index + chan) * 4,
+                                                   1, 1, true, true);
                }
        }
 }
@@ -2892,19 +2722,15 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base)
 static void si_llvm_emit_gs_epilogue(struct lp_build_tgsi_context *bld_base)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef args[2];
 
-       args[0] = lp_build_const_int32(gallivm, SENDMSG_GS_OP_NOP | SENDMSG_GS_DONE);
-       args[1] = LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID);
-       lp_build_intrinsic(gallivm->builder, "llvm.SI.sendmsg",
-                          ctx->voidt, args, 2, 0);
+       ac_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;
@@ -2942,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 = si_llvm_saturate(bld_base, val);
+                               val = ac_build_clamp(&ctx->ac, val);
                                LLVMBuildStore(gallivm->builder, val, addr);
                        }
                }
@@ -2982,7 +2808,7 @@ static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
 
 struct si_ps_exports {
        unsigned num;
-       LLVMValueRef args[10][9];
+       struct ac_export_args args[10];
 };
 
 unsigned si_get_spi_shader_z_format(bool writes_z, bool writes_stencil,
@@ -3010,8 +2836,7 @@ static void si_export_mrt_z(struct lp_build_tgsi_context *bld_base,
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct lp_build_context *base = &bld_base->base;
-       struct lp_build_context *uint = &bld_base->uint_bld;
-       LLVMValueRef args[9];
+       struct ac_export_args args;
        unsigned mask = 0;
        unsigned format = si_get_spi_shader_z_format(depth != NULL,
                                                     stencil != NULL,
@@ -3019,46 +2844,46 @@ static void si_export_mrt_z(struct lp_build_tgsi_context *bld_base,
 
        assert(depth || stencil || samplemask);
 
-       args[1] = uint->one; /* whether the EXEC mask is valid */
-       args[2] = uint->one; /* DONE bit */
+       args.valid_mask = 1; /* whether the EXEC mask is valid */
+       args.done = 1; /* DONE bit */
 
        /* Specify the target we are exporting */
-       args[3] = lp_build_const_int32(base->gallivm, V_008DFC_SQ_EXP_MRTZ);
+       args.target = V_008DFC_SQ_EXP_MRTZ;
 
-       args[4] = uint->zero; /* COMP flag */
-       args[5] = base->undef; /* R, depth */
-       args[6] = base->undef; /* G, stencil test value[0:7], stencil op value[8:15] */
-       args[7] = base->undef; /* B, sample mask */
-       args[8] = base->undef; /* A, alpha to mask */
+       args.compr = 0; /* COMP flag */
+       args.out[0] = base->undef; /* R, depth */
+       args.out[1] = base->undef; /* G, stencil test value[0:7], stencil op value[8:15] */
+       args.out[2] = base->undef; /* B, sample mask */
+       args.out[3] = base->undef; /* A, alpha to mask */
 
        if (format == V_028710_SPI_SHADER_UINT16_ABGR) {
                assert(!depth);
-               args[4] = uint->one; /* COMPR flag */
+               args.compr = 1; /* COMPR flag */
 
                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[5] = bitcast(bld_base, TGSI_TYPE_FLOAT, stencil);
+                       args.out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT, stencil);
                        mask |= 0x3;
                }
                if (samplemask) {
                        /* SampleMask should be in Y[15:0]. */
-                       args[6] = samplemask;
+                       args.out[1] = samplemask;
                        mask |= 0xc;
                }
        } else {
                if (depth) {
-                       args[5] = depth;
+                       args.out[0] = depth;
                        mask |= 0x1;
                }
                if (stencil) {
-                       args[6] = stencil;
+                       args.out[1] = stencil;
                        mask |= 0x2;
                }
                if (samplemask) {
-                       args[7] = samplemask;
+                       args.out[2] = samplemask;
                        mask |= 0x4;
                }
        }
@@ -3071,9 +2896,9 @@ static void si_export_mrt_z(struct lp_build_tgsi_context *bld_base,
                mask |= 0x1;
 
        /* Specify which components to enable */
-       args[0] = lp_build_const_int32(base->gallivm, mask);
+       args.enabled_channels = mask;
 
-       memcpy(exp->args[exp->num++], args, sizeof(args));
+       memcpy(&exp->args[exp->num++], &args, sizeof(args));
 }
 
 static void si_export_mrt_color(struct lp_build_tgsi_context *bld_base,
@@ -3088,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] = si_llvm_saturate(bld_base, color[i]);
+                       color[i] = ac_build_clamp(&ctx->ac, color[i]);
 
        /* Alpha to one */
        if (ctx->shader->key.part.ps.epilog.alpha_to_one)
@@ -3106,40 +2931,40 @@ static void si_export_mrt_color(struct lp_build_tgsi_context *bld_base,
 
        /* If last_cbuf > 0, FS_COLOR0_WRITES_ALL_CBUFS is true. */
        if (ctx->shader->key.part.ps.epilog.last_cbuf > 0) {
-               LLVMValueRef args[8][9];
+               struct ac_export_args args[8];
                int c, last = -1;
 
                /* Get the export arguments, also find out what the last one is. */
                for (c = 0; c <= ctx->shader->key.part.ps.epilog.last_cbuf; c++) {
                        si_llvm_init_export_args(bld_base, color,
-                                                V_008DFC_SQ_EXP_MRT + c, args[c]);
-                       if (args[c][0] != bld_base->uint_bld.zero)
+                                                V_008DFC_SQ_EXP_MRT + c, &args[c]);
+                       if (args[c].enabled_channels)
                                last = c;
                }
 
                /* Emit all exports. */
                for (c = 0; c <= ctx->shader->key.part.ps.epilog.last_cbuf; c++) {
                        if (is_last && last == c) {
-                               args[c][1] = bld_base->uint_bld.one; /* whether the EXEC mask is valid */
-                               args[c][2] = bld_base->uint_bld.one; /* DONE bit */
-                       } else if (args[c][0] == bld_base->uint_bld.zero)
+                               args[c].valid_mask = 1; /* whether the EXEC mask is valid */
+                               args[c].done = 1; /* DONE bit */
+                       } else if (!args[c].enabled_channels)
                                continue; /* unnecessary NULL export */
 
-                       memcpy(exp->args[exp->num++], args[c], sizeof(args[c]));
+                       memcpy(&exp->args[exp->num++], &args[c], sizeof(args[c]));
                }
        } else {
-               LLVMValueRef args[9];
+               struct ac_export_args args;
 
                /* Export */
                si_llvm_init_export_args(bld_base, color, V_008DFC_SQ_EXP_MRT + index,
-                                        args);
+                                        &args);
                if (is_last) {
-                       args[1] = bld_base->uint_bld.one; /* whether the EXEC mask is valid */
-                       args[2] = bld_base->uint_bld.one; /* DONE bit */
-               } else if (args[0] == bld_base->uint_bld.zero)
+                       args.valid_mask = 1; /* whether the EXEC mask is valid */
+                       args.done = 1; /* DONE bit */
+               } else if (!args.enabled_channels)
                        return; /* unnecessary NULL export */
 
-               memcpy(exp->args[exp->num++], args, sizeof(args));
+               memcpy(&exp->args[exp->num++], &args, sizeof(args));
        }
 }
 
@@ -3147,30 +2972,26 @@ static void si_emit_ps_exports(struct si_shader_context *ctx,
                               struct si_ps_exports *exp)
 {
        for (unsigned i = 0; i < exp->num; i++)
-               lp_build_intrinsic(ctx->gallivm.builder,
-                                  "llvm.SI.export", ctx->voidt,
-                                  exp->args[i], 9, 0);
+               ac_build_export(&ctx->ac, &exp->args[i]);
 }
 
 static void si_export_null(struct lp_build_tgsi_context *bld_base)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
        struct lp_build_context *base = &bld_base->base;
-       struct lp_build_context *uint = &bld_base->uint_bld;
-       LLVMValueRef args[9];
+       struct ac_export_args args;
 
-       args[0] = lp_build_const_int32(base->gallivm, 0x0); /* enabled channels */
-       args[1] = uint->one; /* whether the EXEC mask is valid */
-       args[2] = uint->one; /* DONE bit */
-       args[3] = lp_build_const_int32(base->gallivm, V_008DFC_SQ_EXP_NULL);
-       args[4] = uint->zero; /* COMPR flag (0 = 32-bit export) */
-       args[5] = base->undef; /* R */
-       args[6] = base->undef; /* G */
-       args[7] = base->undef; /* B */
-       args[8] = base->undef; /* A */
+       args.enabled_channels = 0x0; /* enabled channels */
+       args.valid_mask = 1; /* whether the EXEC mask is valid */
+       args.done = 1; /* DONE bit */
+       args.target = V_008DFC_SQ_EXP_NULL;
+       args.compr = 0; /* COMPR flag (0 = 32-bit export) */
+       args.out[0] = base->undef; /* R */
+       args.out[1] = base->undef; /* G */
+       args.out[2] = base->undef; /* B */
+       args.out[3] = base->undef; /* A */
 
-       lp_build_intrinsic(base->gallivm->builder, "llvm.SI.export",
-                          ctx->voidt, args, 9, 0);
+       ac_build_export(&ctx->ac, &args);
 }
 
 /**
@@ -3190,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] = {};
@@ -3276,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, "");
        }
@@ -3301,45 +3121,6 @@ static LLVMValueRef get_buffer_size(
        return size;
 }
 
-/**
- * Given the i32 or vNi32 \p type, generate the textual name (e.g. for use with
- * intrinsic names).
- */
-static void build_type_name_for_intr(
-       LLVMTypeRef type,
-       char *buf, unsigned bufsize)
-{
-       LLVMTypeRef elem_type = type;
-
-       assert(bufsize >= 8);
-
-       if (LLVMGetTypeKind(type) == LLVMVectorTypeKind) {
-               int ret = snprintf(buf, bufsize, "v%u",
-                                       LLVMGetVectorSize(type));
-               if (ret < 0) {
-                       char *type_name = LLVMPrintTypeToString(type);
-                       fprintf(stderr, "Error building type name for: %s\n",
-                               type_name);
-                       return;
-               }
-               elem_type = LLVMGetElementType(type);
-               buf += ret;
-               bufsize -= ret;
-       }
-       switch (LLVMGetTypeKind(elem_type)) {
-       default: break;
-       case LLVMIntegerTypeKind:
-               snprintf(buf, bufsize, "i%d", LLVMGetIntTypeWidth(elem_type));
-               break;
-       case LLVMFloatTypeKind:
-               snprintf(buf, bufsize, "f32");
-               break;
-       case LLVMDoubleTypeKind:
-               snprintf(buf, bufsize, "f64");
-               break;
-       }
-}
-
 static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
                                struct lp_build_tgsi_context *bld_base,
                                struct lp_build_emit_data *emit_data);
@@ -3347,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
@@ -3368,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);
@@ -3399,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)
@@ -3414,7 +3241,7 @@ shader_buffer_fetch_rsrc(struct si_shader_context *ctx,
                                                   reg->Register.Index,
                                                   SI_NUM_SHADER_BUFFERS);
 
-       return build_indexed_load_const(ctx, rsrc_ptr, index);
+       return ac_build_indexed_load_const(&ctx->ac, rsrc_ptr, index);
 }
 
 static bool tgsi_is_array_sampler(unsigned target)
@@ -3472,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.
  */
@@ -3485,18 +3330,19 @@ 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);
 
        if (!image->Register.Indirect) {
                const struct tgsi_shader_info *info = bld_base->info;
+               unsigned images_writemask = info->images_store |
+                                           info->images_atomic;
 
                index = LLVMConstInt(ctx->i32, image->Register.Index, 0);
 
-               if (info->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:
@@ -3513,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 = build_indexed_load_const(ctx, rsrc_ptr, index);
-               return;
-       }
-
-       tmp = build_indexed_load_const(ctx, 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(
@@ -3537,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);
@@ -3547,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];
 
@@ -3639,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;
@@ -3654,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;
@@ -3666,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);
@@ -3678,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;
@@ -3707,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,
@@ -3735,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);
-               derived_ptr = LLVMBuildGEP(builder, ptr, &index, 1, "");
-               channels[chan] = LLVMBuildLoad(builder, derived_ptr, "");
+
+               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);
+}
+
+/**
+ * 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;
+               }
+
+               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;
+               }
        }
-       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];
-
-       build_type_name_for_intr(coords_type, coords_type_name,
-                           sizeof(coords_type_name));
 
-       if (HAVE_LLVM <= 0x0309) {
-               snprintf(out_name, out_len, "%s.%s", base_name, coords_type_name);
+       /* 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];
-
-               build_type_name_for_intr(data_type, data_type_name,
-                                       sizeof(data_type_name));
-               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(
@@ -3789,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);
@@ -3802,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;
        }
 
@@ -3812,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));
        }
 }
 
@@ -3833,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;
@@ -3860,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;
@@ -3881,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);
@@ -3896,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;
@@ -3930,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";
                }
 
@@ -3954,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;
@@ -3962,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));
        }
 }
 
@@ -3972,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);
        }
@@ -3997,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);
@@ -4011,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;
        }
 
@@ -4020,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));
        }
 }
 
@@ -4040,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
@@ -4070,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;
@@ -4083,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;
@@ -4168,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];
@@ -4192,17 +4107,77 @@ static void atomic_emit(
                else
                        coords = emit_data->args[1];
 
-               build_type_name_for_intr(LLVMTypeOf(coords), coords_type, sizeof(coords_type));
+               ac_build_type_name_for_intr(LLVMTypeOf(coords), coords_type, sizeof(coords_type));
                snprintf(intrinsic_name, sizeof(intrinsic_name),
                         "llvm.amdgcn.image.atomic.%s.%s",
                         action->intr_name, coords_type);
        }
 
        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(
@@ -4210,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];
 
@@ -4224,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);
        }
 }
 
@@ -4245,82 +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);
+               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;
-
-       /* 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 {
@@ -4333,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;
@@ -4348,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]. */
@@ -4366,16 +4288,7 @@ static LLVMValueRef load_sampler_desc_custom(struct si_shader_context *ctx,
                break;
        }
 
-       return build_indexed_load_const(ctx, 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);
+       return ac_build_indexed_load_const(&ctx->ac, list, index);
 }
 
 /* Disable anisotropic filtering if BASE_LEVEL == LAST_LEVEL.
@@ -4401,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(
@@ -4413,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;
@@ -4446,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);
                }
        }
@@ -4476,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);
 }
 
@@ -4484,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) {
@@ -4493,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);
-
-       /* 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(
@@ -4519,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;
@@ -4540,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;
@@ -4563,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].
                 */
@@ -4575,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], "");
@@ -4611,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 = si_llvm_saturate(bld_base, 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:
@@ -4632,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:
@@ -4640,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:
@@ -4647,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 ||
@@ -4680,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];
@@ -4712,25 +4655,20 @@ 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];
-               unsigned txf_count = count;
+               /* We only need .xy for non-arrays, and .xyz for arrays. */
+               unsigned txf_count = target == TGSI_TEXTURE_2D_MSAA ? 2 : 3;
                struct tgsi_full_instruction inst = {};
 
                memcpy(txf_address, address, sizeof(txf_address));
 
-               if (target == TGSI_TEXTURE_2D_MSAA) {
-                       txf_address[2] = bld_base->uint_bld.zero;
-               }
-               txf_address[3] = bld_base->uint_bld.zero;
-
-               /* Read FMASK using TXF. */
-               inst.Instruction.Opcode = TGSI_OPCODE_TXF;
+               /* 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);
@@ -4743,9 +4681,9 @@ static void tex_fetch_args(
                LLVMValueRef fmask =
                        LLVMBuildExtractElement(gallivm->builder,
                                                txf_emit_data.output[0],
-                                               uint_bld->zero, "");
+                                               ctx->i32_0, "");
 
-               unsigned sample_chan = target == TGSI_TEXTURE_2D_MSAA ? 2 : 3;
+               unsigned sample_chan = txf_count; /* the sample index is last */
 
                LLVMValueRef sample_index4 =
                        LLVMBuildMul(gallivm->builder, address[sample_chan], four, "");
@@ -4765,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] =
@@ -4777,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;
@@ -4839,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);
 }
 
@@ -4853,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);
 
@@ -4905,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);
+       args->addr = coord;
 }
 
 static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
@@ -4917,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);
+               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. */
-       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.
@@ -4996,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);
+       emit_data->output[emit_data->chan] =
+               ac_build_image_opcode(&ctx->ac, &args);
 }
 
 static void si_llvm_emit_txqs(
@@ -5019,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;
@@ -5031,107 +4959,43 @@ 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;
 }
 
-/*
- * SI implements derivatives using the local data store (LDS)
- * All writes to the LDS happen in all executing threads at
- * the same time. TID is the Thread ID for the current
- * thread and is a value between 0 and 63, representing
- * the thread's position in the wavefront.
- *
- * For the pixel shader threads are grouped into quads of four pixels.
- * The TIDs of the pixels of a quad are:
- *
- *  +------+------+
- *  |4n + 0|4n + 1|
- *  +------+------+
- *  |4n + 2|4n + 3|
- *  +------+------+
- *
- * So, masking the TID with 0xfffffffc yields the TID of the top left pixel
- * of the quad, masking with 0xfffffffd yields the TID of the top pixel of
- * the current pixel's column, and masking with 0xfffffffe yields the TID
- * of the left pixel of the current pixel's row.
- *
- * Adding 1 yields the TID of the pixel to the right of the left pixel, and
- * adding 2 yields the TID of the pixel below the top pixel.
- */
-/* masks for thread ID. */
-#define TID_MASK_TOP_LEFT 0xfffffffc
-#define TID_MASK_TOP      0xfffffffd
-#define TID_MASK_LEFT     0xfffffffe
-
 static void si_llvm_emit_ddxy(
        const struct lp_build_tgsi_action *action,
        struct lp_build_tgsi_context *bld_base,
        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 thread_id, tl, trbl, tl_tid, trbl_tid, val, args[2];
+       LLVMValueRef val;
        int idx;
        unsigned mask;
 
-       thread_id = get_thread_id(ctx);
-
        if (opcode == TGSI_OPCODE_DDX_FINE)
-               mask = TID_MASK_LEFT;
+               mask = AC_TID_MASK_LEFT;
        else if (opcode == TGSI_OPCODE_DDY_FINE)
-               mask = TID_MASK_TOP;
+               mask = AC_TID_MASK_TOP;
        else
-               mask = TID_MASK_TOP_LEFT;
-
-       tl_tid = LLVMBuildAnd(gallivm->builder, thread_id,
-                               lp_build_const_int32(gallivm, mask), "");
+               mask = AC_TID_MASK_TOP_LEFT;
 
        /* for DDX we want to next X pixel, DDY next Y pixel. */
        idx = (opcode == TGSI_OPCODE_DDX || opcode == TGSI_OPCODE_DDX_FINE) ? 1 : 2;
-       trbl_tid = LLVMBuildAdd(gallivm->builder, tl_tid,
-                                 lp_build_const_int32(gallivm, idx), "");
 
        val = LLVMBuildBitCast(gallivm->builder, emit_data->args[0], ctx->i32, "");
-
-       if (ctx->screen->has_ds_bpermute) {
-               args[0] = LLVMBuildMul(gallivm->builder, tl_tid,
-                                      lp_build_const_int32(gallivm, 4), "");
-               args[1] = val;
-               tl = lp_build_intrinsic(gallivm->builder,
-                                       "llvm.amdgcn.ds.bpermute", ctx->i32,
-                                       args, 2, LP_FUNC_ATTR_READNONE);
-
-               args[0] = LLVMBuildMul(gallivm->builder, trbl_tid,
-                                      lp_build_const_int32(gallivm, 4), "");
-               trbl = lp_build_intrinsic(gallivm->builder,
-                                         "llvm.amdgcn.ds.bpermute", ctx->i32,
-                                         args, 2, LP_FUNC_ATTR_READNONE);
-       } else {
-               LLVMValueRef store_ptr, load_ptr0, load_ptr1;
-
-               store_ptr = build_gep0(ctx, ctx->lds, thread_id);
-               load_ptr0 = build_gep0(ctx, ctx->lds, tl_tid);
-               load_ptr1 = build_gep0(ctx, ctx->lds, trbl_tid);
-
-               LLVMBuildStore(gallivm->builder, val, store_ptr);
-               tl = LLVMBuildLoad(gallivm->builder, load_ptr0, "");
-               trbl = LLVMBuildLoad(gallivm->builder, load_ptr1, "");
-       }
-
-       tl = LLVMBuildBitCast(gallivm->builder, tl, ctx->f32, "");
-       trbl = LLVMBuildBitCast(gallivm->builder, trbl, ctx->f32, "");
-
-       emit_data->output[emit_data->chan] =
-               LLVMBuildFSub(gallivm->builder, trbl, tl, "");
+       val = ac_build_ddxy(&ctx->ac, ctx->screen->has_ds_bpermute,
+                           mask, idx, ctx->lds, val);
+       emit_data->output[emit_data->chan] = val;
 }
 
 /*
@@ -5144,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;
 
@@ -5163,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) {
@@ -5178,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.
@@ -5191,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;
        }
@@ -5208,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;
@@ -5237,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) {
@@ -5253,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,
@@ -5274,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++) {
@@ -5282,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, "");
-                       emit_data->output[chan] = build_fs_interp(bld_base,
+                               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] = build_fs_interp_mov(bld_base,
-                               lp_build_const_int32(gallivm, 2), /* P0 */
+                       emit_data->output[chan] = ac_build_fs_interp_mov(&ctx->ac,
+                               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)
 {
@@ -5327,13 +5328,12 @@ 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);
        LLVMValueRef gs_next_vertex;
        LLVMValueRef can_emit, kill;
-       LLVMValueRef args[2];
        unsigned chan, offset;
        int i;
        unsigned stream;
@@ -5354,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, 0);
+               ac_build_kill(&ctx->ac, kill);
        } else {
                lp_build_if(&if_state, gallivm, can_emit);
        }
@@ -5380,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);
@@ -5389,27 +5388,22 @@ static void si_llvm_emit_vertex(
 
                        out_val = LLVMBuildBitCast(gallivm->builder, out_val, ctx->i32, "");
 
-                       build_tbuffer_store(ctx,
-                                           ctx->gsvs_ring[stream],
-                                           out_val, 1,
-                                           voffset, soffset, 0,
-                                           V_008F0C_BUF_DATA_FORMAT_32,
-                                           V_008F0C_BUF_NUM_FORMAT_UINT,
-                                           1, 0, 1, 1, 0);
+                       ac_build_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 */
-       args[0] = lp_build_const_int32(gallivm, SENDMSG_GS_OP_EMIT | SENDMSG_GS | (stream << 8));
-       args[1] = LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID);
-       lp_build_intrinsic(gallivm->builder, "llvm.SI.sendmsg",
-                          ctx->voidt, args, 2, 0);
-
+       ac_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);
 }
@@ -5421,16 +5415,12 @@ static void si_llvm_emit_primitive(
        struct lp_build_emit_data *emit_data)
 {
        struct si_shader_context *ctx = si_shader_context(bld_base);
-       struct gallivm_state *gallivm = bld_base->base.gallivm;
-       LLVMValueRef args[2];
        unsigned stream;
 
        /* Signal primitive cut */
        stream = si_llvm_get_stream(bld_base, emit_data);
-       args[0] = lp_build_const_int32(gallivm, SENDMSG_GS_OP_CUT | SENDMSG_GS | (stream << 8));
-       args[1] = LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID);
-       lp_build_intrinsic(gallivm->builder, "llvm.SI.sendmsg",
-                          ctx->voidt, args, 2, 0);
+       ac_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,
@@ -5438,12 +5428,15 @@ 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;
 
-       /* The real barrier instruction isn’t needed, because an entire patch
+       /* SI only (thanks to a hw bug workaround):
+        * The real barrier instruction isn’t needed, because an entire patch
         * always fits into a single wave.
         */
-       if (ctx->type == PIPE_SHADER_TESS_CTRL) {
+       if (HAVE_LLVM >= 0x0309 &&
+           ctx->screen->b.chip_class == SI &&
+           ctx->type == PIPE_SHADER_TESS_CTRL) {
                emit_waitcnt(ctx, LGKM_CNT & VM_CNT);
                return;
        }
@@ -5451,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 = {
@@ -5489,11 +5482,16 @@ 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);
        }
 
+       LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
+                                          "no-signed-zeros-fp-math",
+                                          "true");
+
        if (ctx->screen->b.debug_flags & DBG_UNSAFE_MATH) {
                /* These were copied from some LLVM test. */
                LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
@@ -5511,20 +5509,6 @@ static void si_create_function(struct si_shader_context *ctx,
        }
 }
 
-static void create_meta_data(struct si_shader_context *ctx)
-{
-       struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
-
-       ctx->invariant_load_md_kind = LLVMGetMDKindIDInContext(gallivm->context,
-                                                              "invariant.load", 14);
-       ctx->range_md_kind = LLVMGetMDKindIDInContext(gallivm->context,
-                                                    "range", 5);
-       ctx->uniform_md_kind = LLVMGetMDKindIDInContext(gallivm->context,
-                                                       "amdgpu.uniform", 14);
-
-       ctx->empty_md = LLVMMDNodeInContext(gallivm->context, NULL, 0);
-}
-
 static void declare_streamout_params(struct si_shader_context *ctx,
                                     struct pipe_stream_output_info *so,
                                     LLVMTypeRef *params, LLVMTypeRef i32,
@@ -5537,7 +5521,7 @@ static void declare_streamout_params(struct si_shader_context *ctx,
                if (ctx->type != PIPE_SHADER_TESS_EVAL)
                        params[ctx->param_streamout_config = (*num_params)++] = i32;
                else
-                       ctx->param_streamout_config = ctx->param_tess_offchip;
+                       ctx->param_streamout_config = *num_params - 1;
 
                params[ctx->param_streamout_write_index = (*num_params)++] = i32;
        }
@@ -5576,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");
 }
@@ -5605,9 +5587,9 @@ 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_NUM_VERTEX_BUFFERS], v3i32;
+       LLVMTypeRef params[SI_NUM_PARAMS + SI_MAX_ATTRIBS], v3i32;
        LLVMTypeRef returns[16+32*4];
        unsigned i, last_sgpr, num_params, num_return_sgprs;
        unsigned num_returns = 0;
@@ -5623,7 +5605,7 @@ static void create_function(struct si_shader_context *ctx)
 
        switch (ctx->type) {
        case PIPE_SHADER_VERTEX:
-               params[SI_PARAM_VERTEX_BUFFERS] = const_array(ctx->v16i8, SI_NUM_VERTEX_BUFFERS);
+               params[SI_PARAM_VERTEX_BUFFERS] = const_array(ctx->v16i8, SI_MAX_ATTRIBS);
                params[SI_PARAM_BASE_VERTEX] = ctx->i32;
                params[SI_PARAM_START_INSTANCE] = ctx->i32;
                params[SI_PARAM_DRAWID] = ctx->i32;
@@ -5701,10 +5683,10 @@ static void create_function(struct si_shader_context *ctx)
 
                if (shader->key.as_es) {
                        params[ctx->param_oc_lds = num_params++] = ctx->i32;
-                       params[ctx->param_tess_offchip = num_params++] = ctx->i32;
+                       params[num_params++] = ctx->i32;
                        params[ctx->param_es2gs_offset = num_params++] = ctx->i32;
                } else {
-                       params[ctx->param_tess_offchip = num_params++] = ctx->i32;
+                       params[num_params++] = ctx->i32;
                        declare_streamout_params(ctx, &shader->selector->so,
                                                 params, ctx->i32, &num_params);
                        params[ctx->param_oc_lds = num_params++] = ctx->i32;
@@ -5860,8 +5842,7 @@ static void create_function(struct si_shader_context *ctx)
                                                    LOCAL_ADDR_SPACE);
 
        if ((ctx->type == PIPE_SHADER_VERTEX && shader->key.as_ls) ||
-           ctx->type == PIPE_SHADER_TESS_CTRL ||
-           ctx->type == PIPE_SHADER_TESS_EVAL)
+           ctx->type == PIPE_SHADER_TESS_CTRL)
                declare_tess_lds(ctx);
 }
 
@@ -5871,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,
@@ -5885,24 +5866,23 @@ 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 =
-                       build_indexed_load_const(ctx, buf_ptr, offset);
+                       ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
        }
 
        if (ctx->shader->is_gs_copy_shader) {
-               LLVMValueRef offset = lp_build_const_int32(gallivm, SI_RING_GSVS);
+               LLVMValueRef offset = LLVMConstInt(ctx->i32, SI_RING_GSVS, 0);
 
                ctx->gsvs_ring[0] =
-                       build_indexed_load_const(ctx, buf_ptr, offset);
+                       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 = build_indexed_load_const(ctx, buf_ptr, offset);
+               base_ring = ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
 
                /* The conceptual layout of the GSVS ring is
                 *   v0c0 .. vLv0 v0c1 .. vLc1 ..
@@ -5933,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), "");
@@ -5974,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];
 
@@ -5987,8 +5966,8 @@ 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);
-       desc = build_indexed_load_const(ctx, param_rw_buffers, slot);
+       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. */
        offset = LLVMBuildMul(builder, address[1],
@@ -6001,23 +5980,23 @@ 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, 0);
+       ac_build_kill(&ctx->ac, bit);
 }
 
-void si_shader_binary_read_config(struct radeon_shader_binary *binary,
+void si_shader_binary_read_config(struct ac_shader_binary *binary,
                                  struct si_shader_config *conf,
                                  unsigned symbol_offset)
 {
        unsigned i;
        const unsigned char *config =
-               radeon_shader_binary_config_start(binary, symbol_offset);
+               ac_shader_binary_config_start(binary, symbol_offset);
        bool really_needs_scratch = false;
 
        /* LLVM adds SGPR spills to the scratch size.
         * Find out if we really need the scratch buffer.
         */
        for (i = 0; i < binary->reloc_count; i++) {
-               const struct radeon_shader_reloc *reloc = &binary->relocs[i];
+               const struct ac_shader_reloc *reloc = &binary->relocs[i];
 
                if (!strcmp(scratch_rsrc_dword0_symbol, reloc->name) ||
                    !strcmp(scratch_rsrc_dword1_symbol, reloc->name)) {
@@ -6107,7 +6086,7 @@ void si_shader_apply_scratch_relocs(struct si_context *sctx,
                        S_008F04_STRIDE(config->scratch_bytes_per_wave / 64);
 
        for (i = 0 ; i < shader->binary.reloc_count; i++) {
-               const struct radeon_shader_reloc *reloc =
+               const struct ac_shader_reloc *reloc =
                                        &shader->binary.relocs[i];
                if (!strcmp(scratch_rsrc_dword0_symbol, reloc->name)) {
                        util_memcpy_cpu_to_le32(shader->binary.code + reloc->offset,
@@ -6132,11 +6111,11 @@ static unsigned si_get_shader_binary_size(struct si_shader *shader)
 
 int si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader)
 {
-       const struct radeon_shader_binary *prolog =
+       const struct ac_shader_binary *prolog =
                shader->prolog ? &shader->prolog->binary : NULL;
-       const struct radeon_shader_binary *epilog =
+       const struct ac_shader_binary *epilog =
                shader->epilog ? &shader->epilog->binary : NULL;
-       const struct radeon_shader_binary *mainb = &shader->binary;
+       const struct ac_shader_binary *mainb = &shader->binary;
        unsigned bo_size = si_get_shader_binary_size(shader) +
                           (!epilog ? mainb->rodata_size : 0);
        unsigned char *ptr;
@@ -6145,10 +6124,17 @@ 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,
-                                       PIPE_USAGE_IMMUTABLE, bo_size);
+                                       PIPE_USAGE_IMMUTABLE,
+                                       align(bo_size, SI_CPDMA_ALIGNMENT));
        if (!shader->bo)
                return -ENOMEM;
 
@@ -6173,7 +6159,7 @@ int si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader)
        return 0;
 }
 
-static void si_shader_dump_disassembly(const struct radeon_shader_binary *binary,
+static void si_shader_dump_disassembly(const struct ac_shader_binary *binary,
                                       struct pipe_debug_callback *debug,
                                       const char *name, FILE *file)
 {
@@ -6314,8 +6300,7 @@ static void si_shader_dump_stats(struct si_screen *sscreen,
                           conf->spilled_vgprs, conf->private_mem_vgprs);
 }
 
-static const char *si_get_shader_name(struct si_shader *shader,
-                                     unsigned processor)
+const char *si_get_shader_name(struct si_shader *shader, unsigned processor)
 {
        switch (processor) {
        case PIPE_SHADER_VERTEX:
@@ -6382,7 +6367,7 @@ void si_shader_dump(struct si_screen *sscreen, struct si_shader *shader,
 }
 
 int si_compile_llvm(struct si_screen *sscreen,
-                   struct radeon_shader_binary *binary,
+                   struct ac_shader_binary *binary,
                    struct si_shader_config *conf,
                    LLVMTargetMachineRef tm,
                    LLVMModuleRef mod,
@@ -6398,7 +6383,7 @@ int si_compile_llvm(struct si_screen *sscreen,
 
                if (!(sscreen->b.debug_flags & (DBG_NO_IR | DBG_PREOPT_IR))) {
                        fprintf(stderr, "%s LLVM IR:\n\n", name);
-                       LLVMDumpModule(mod);
+                       ac_dump_module(mod);
                        fprintf(stderr, "\n");
                }
        }
@@ -6474,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]));
@@ -6497,21 +6481,12 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
 
        builder = gallivm->builder;
 
-       create_meta_data(&ctx);
        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;
@@ -6519,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) {
@@ -6549,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 */
@@ -6562,18 +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),
-                                                ctx.f32, "");
+                                       ac_build_buffer_load(&ctx.ac,
+                                                            ctx.gsvs_ring[0], 1,
+                                                            ctx.i32_0, voffset,
+                                                            soffset, 0, 1, 1, true);
                        }
                }
 
@@ -6597,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))
-               LLVMDumpModule(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) {
@@ -6643,7 +6615,11 @@ static void si_dump_shader_key(unsigned shader, struct si_shader_key *key,
                fprintf(f, "  part.vs.epilog.export_prim_id = %u\n", key->part.vs.epilog.export_prim_id);
                fprintf(f, "  as_es = %u\n", key->as_es);
                fprintf(f, "  as_ls = %u\n", key->as_ls);
-               fprintf(f, "  mono.vs.fix_fetch = 0x%"PRIx64"\n", key->mono.vs.fix_fetch);
+
+               fprintf(f, "  mono.vs.fix_fetch = {");
+               for (i = 0; i < SI_MAX_ATTRIBS; i++)
+                       fprintf(f, !i ? "%u" : ", %u", key->mono.vs.fix_fetch[i]);
+               fprintf(f, "}\n");
                break;
 
        case PIPE_SHADER_TESS_CTRL:
@@ -6675,6 +6651,7 @@ static void si_dump_shader_key(unsigned shader, struct si_shader_key *key,
                fprintf(f, "  part.ps.prolog.bc_optimize_for_linear = %u\n", key->part.ps.prolog.bc_optimize_for_linear);
                fprintf(f, "  part.ps.epilog.spi_shader_col_format = 0x%x\n", key->part.ps.epilog.spi_shader_col_format);
                fprintf(f, "  part.ps.epilog.color_is_int8 = 0x%X\n", key->part.ps.epilog.color_is_int8);
+               fprintf(f, "  part.ps.epilog.color_is_int10 = 0x%X\n", key->part.ps.epilog.color_is_int10);
                fprintf(f, "  part.ps.epilog.last_cbuf = %u\n", key->part.ps.epilog.last_cbuf);
                fprintf(f, "  part.ps.epilog.alpha_func = %u\n", key->part.ps.epilog.alpha_func);
                fprintf(f, "  part.ps.epilog.alpha_to_one = %u\n", key->part.ps.epilog.alpha_to_one);
@@ -6716,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;
@@ -6762,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)
@@ -6783,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)) {
@@ -6867,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)
@@ -6912,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;
@@ -6995,7 +6991,6 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx,
                return false;
        }
 
-       create_meta_data(ctx);
        create_function(ctx);
        preload_ring_buffers(ctx);
 
@@ -7003,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, "");
                }
        }
@@ -7483,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;
 
@@ -7503,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)) {
@@ -7596,12 +7589,12 @@ 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 &&
            r600_can_dump_shader(&sscreen->b, ctx.type))
-               LLVMDumpModule(mod);
+               ac_dump_module(mod);
 
        si_llvm_finalize_module(&ctx,
                                    r600_extra_shader_checks(&sscreen->b, ctx.type));
@@ -7728,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;
                }
        }
@@ -7787,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;
 }
 
@@ -7920,24 +7913,21 @@ static void si_build_vs_epilog_function(struct si_shader_context *ctx,
        /* Emit exports. */
        if (key->vs_epilog.states.export_prim_id) {
                struct lp_build_context *base = &bld_base->base;
-               struct lp_build_context *uint = &bld_base->uint_bld;
-               LLVMValueRef args[9];
-
-               args[0] = lp_build_const_int32(base->gallivm, 0x0); /* enabled channels */
-               args[1] = uint->zero; /* whether the EXEC mask is valid */
-               args[2] = uint->zero; /* DONE bit */
-               args[3] = lp_build_const_int32(base->gallivm, V_008DFC_SQ_EXP_PARAM +
-                                              key->vs_epilog.prim_id_param_offset);
-               args[4] = uint->zero; /* COMPR flag (0 = 32-bit export) */
-               args[5] = LLVMGetParam(ctx->main_fn,
+               struct ac_export_args args;
+
+               args.enabled_channels = 0x1; /* enabled channels */
+               args.valid_mask = 0; /* whether the EXEC mask is valid */
+               args.done = 0; /* DONE bit */
+               args.target = V_008DFC_SQ_EXP_PARAM +
+                             key->vs_epilog.prim_id_param_offset;
+               args.compr = 0; /* COMPR flag (0 = 32-bit export) */
+               args.out[0] = LLVMGetParam(ctx->main_fn,
                                       VS_EPILOG_PRIMID_LOC); /* X */
-               args[6] = base->undef; /* Y */
-               args[7] = base->undef; /* Z */
-               args[8] = base->undef; /* W */
+               args.out[1] = base->undef; /* Y */
+               args.out[2] = base->undef; /* Z */
+               args.out[3] = base->undef; /* W */
 
-               lp_build_intrinsic(base->gallivm->builder, "llvm.SI.export",
-                                  LLVMVoidTypeInContext(base->gallivm->context),
-                                  args, 9, 0);
+               ac_build_export(&ctx->ac, &args);
        }
 
        LLVMBuildRetVoid(gallivm->builder);
@@ -8580,7 +8570,7 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
                     struct pipe_debug_callback *debug)
 {
        struct si_shader_selector *sel = shader->selector;
-       struct si_shader *mainp = sel->main_shader_part;
+       struct si_shader *mainp = *si_get_main_shader_part(sel, &shader->key);
        int r;
 
        /* LS, ES, VS are compiled on demand if the main part hasn't been