X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=src%2Fgallium%2Fdrivers%2Fradeonsi%2Fsi_shader.c;h=934639822d0b011d21c5ea447a0066c2b17240ed;hb=0c6c6810bd2a1e98c0ae07d6db981b8d4f79856f;hp=c18915488e53f8cd706ab8ad159b4b6f116c60d6;hpb=5f1cddde78aa93ea1272c50a93e479cb76144af7;p=mesa.git

diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index c18915488e5..934639822d0 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -22,62 +22,34 @@
  * USE OR OTHER DEALINGS IN THE SOFTWARE.
  */
 
-#include "gallivm/lp_bld_const.h"
-#include "gallivm/lp_bld_gather.h"
-#include "gallivm/lp_bld_intr.h"
-#include "gallivm/lp_bld_logic.h"
-#include "gallivm/lp_bld_arit.h"
-#include "gallivm/lp_bld_flow.h"
-#include "gallivm/lp_bld_misc.h"
 #include "util/u_memory.h"
 #include "util/u_string.h"
 #include "tgsi/tgsi_build.h"
+#include "tgsi/tgsi_strings.h"
 #include "tgsi/tgsi_util.h"
 #include "tgsi/tgsi_dump.h"
+#include "tgsi/tgsi_from_mesa.h"
 
 #include "ac_binary.h"
-#include "ac_llvm_util.h"
 #include "ac_exp_param.h"
 #include "ac_shader_util.h"
+#include "ac_rtld.h"
+#include "ac_llvm_util.h"
 #include "si_shader_internal.h"
 #include "si_pipe.h"
 #include "sid.h"
 
 #include "compiler/nir/nir.h"
 
-static const char *scratch_rsrc_dword0_symbol =
+static const char scratch_rsrc_dword0_symbol[] =
 	"SCRATCH_RSRC_DWORD0";
 
-static const char *scratch_rsrc_dword1_symbol =
+static const char scratch_rsrc_dword1_symbol[] =
 	"SCRATCH_RSRC_DWORD1";
 
-struct si_shader_output_values
-{
-	LLVMValueRef values[4];
-	unsigned semantic_name;
-	unsigned semantic_index;
-	ubyte vertex_stream[4];
-};
-
-/**
- * Used to collect types and other info about arguments of the LLVM function
- * before the function is created.
- */
-struct si_function_info {
-	LLVMTypeRef types[100];
-	LLVMValueRef *assign[100];
-	unsigned num_sgpr_params;
-	unsigned num_params;
-};
-
-enum si_arg_regfile {
-	ARG_SGPR,
-	ARG_VGPR
-};
-
 static void si_init_shader_ctx(struct si_shader_context *ctx,
 			       struct si_screen *sscreen,
-			       LLVMTargetMachineRef tm);
+			       struct ac_llvm_compiler *compiler);
 
 static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
 				 struct lp_build_tgsi_context *bld_base,
@@ -94,6 +66,8 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx,
 					union si_shader_part_key *key);
 static void si_build_ps_epilog_function(struct si_shader_context *ctx,
 					union si_shader_part_key *key);
+static void si_fix_resource_usage(struct si_screen *sscreen,
+				  struct si_shader *shader);
 
 /* Ideally pass the sample mask input to the PS epilog as v14, which
  * is its usual location, so that the shader doesn't have to add v_mov.
@@ -109,24 +83,24 @@ static bool llvm_type_is_64bit(struct si_shader_context *ctx,
 	return false;
 }
 
-static bool is_merged_shader(struct si_shader *shader)
+static bool is_merged_shader(struct si_shader_context *ctx)
 {
-	if (shader->selector->screen->info.chip_class <= VI)
+	if (ctx->screen->info.chip_class <= GFX8)
 		return false;
 
-	return shader->key.as_ls ||
-	       shader->key.as_es ||
-	       shader->selector->type == PIPE_SHADER_TESS_CTRL ||
-	       shader->selector->type == PIPE_SHADER_GEOMETRY;
+	return ctx->shader->key.as_ls ||
+	       ctx->shader->key.as_es ||
+	       ctx->type == PIPE_SHADER_TESS_CTRL ||
+	       ctx->type == PIPE_SHADER_GEOMETRY;
 }
 
-static void si_init_function_info(struct si_function_info *fninfo)
+void si_init_function_info(struct si_function_info *fninfo)
 {
 	fninfo->num_params = 0;
 	fninfo->num_sgpr_params = 0;
 }
 
-static unsigned add_arg_assign(struct si_function_info *fninfo,
+unsigned add_arg_assign(struct si_function_info *fninfo,
 			enum si_arg_regfile regfile, LLVMTypeRef type,
 			LLVMValueRef *assign)
 {
@@ -191,7 +165,8 @@ unsigned si_shader_io_get_unique_index_patch(unsigned semantic_name, unsigned in
  * less than 64, so that a 64-bit bitmask of used inputs or outputs can be
  * calculated.
  */
-unsigned si_shader_io_get_unique_index(unsigned semantic_name, unsigned index)
+unsigned si_shader_io_get_unique_index(unsigned semantic_name, unsigned index,
+				       unsigned is_varying)
 {
 	switch (semantic_name) {
 	case TGSI_SEMANTIC_POSITION:
@@ -220,15 +195,24 @@ unsigned si_shader_io_get_unique_index(unsigned semantic_name, unsigned index)
 		return SI_MAX_IO_GENERIC + 6;
 	case TGSI_SEMANTIC_PRIMID:
 		return SI_MAX_IO_GENERIC + 7;
-	case TGSI_SEMANTIC_COLOR: /* these alias */
-	case TGSI_SEMANTIC_BCOLOR:
+	case TGSI_SEMANTIC_COLOR:
 		assert(index < 2);
 		return SI_MAX_IO_GENERIC + 8 + index;
+	case TGSI_SEMANTIC_BCOLOR:
+		assert(index < 2);
+		/* If it's a varying, COLOR and BCOLOR alias. */
+		if (is_varying)
+			return SI_MAX_IO_GENERIC + 8 + index;
+		else
+			return SI_MAX_IO_GENERIC + 10 + index;
 	case TGSI_SEMANTIC_TEXCOORD:
 		assert(index < 8);
-		assert(SI_MAX_IO_GENERIC + 10 + index < 64);
-		return SI_MAX_IO_GENERIC + 10 + index;
+		STATIC_ASSERT(SI_MAX_IO_GENERIC + 12 + 8 <= 63);
+		return SI_MAX_IO_GENERIC + 12 + index;
+	case TGSI_SEMANTIC_CLIPVERTEX:
+		return 63;
 	default:
+		fprintf(stderr, "invalid semantic name = %u\n", semantic_name);
 		assert(!"invalid semantic name");
 		return 0;
 	}
@@ -343,21 +327,21 @@ static LLVMValueRef get_tcs_out_patch_stride(struct si_shader_context *ctx)
 static LLVMValueRef
 get_tcs_out_patch0_offset(struct si_shader_context *ctx)
 {
-	return lp_build_mul_imm(&ctx->bld_base.uint_bld,
-				si_unpack_param(ctx,
-					     ctx->param_tcs_out_lds_offsets,
-					     0, 16),
-				4);
+	return LLVMBuildMul(ctx->ac.builder,
+			    si_unpack_param(ctx,
+					    ctx->param_tcs_out_lds_offsets,
+					    0, 16),
+			    LLVMConstInt(ctx->i32, 4, 0), "");
 }
 
 static LLVMValueRef
 get_tcs_out_patch0_patch_data_offset(struct si_shader_context *ctx)
 {
-	return lp_build_mul_imm(&ctx->bld_base.uint_bld,
-				si_unpack_param(ctx,
-					     ctx->param_tcs_out_lds_offsets,
-					     16, 16),
-				4);
+	return LLVMBuildMul(ctx->ac.builder,
+			    si_unpack_param(ctx,
+					    ctx->param_tcs_out_lds_offsets,
+					    16, 16),
+			    LLVMConstInt(ctx->i32, 4, 0), "");
 }
 
 static LLVMValueRef
@@ -376,10 +360,7 @@ get_tcs_out_current_patch_offset(struct si_shader_context *ctx)
 	LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
 	LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
 
-	return LLVMBuildAdd(ctx->ac.builder, patch0_offset,
-			    LLVMBuildMul(ctx->ac.builder, patch_stride,
-					 rel_patch_id, ""),
-			    "");
+	return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id, patch0_offset);
 }
 
 static LLVMValueRef
@@ -390,10 +371,7 @@ get_tcs_out_current_patch_data_offset(struct si_shader_context *ctx)
 	LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
 	LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
 
-	return LLVMBuildAdd(ctx->ac.builder, patch0_patch_data_offset,
-			    LLVMBuildMul(ctx->ac.builder, patch_stride,
-					 rel_patch_id, ""),
-			    "");
+	return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id, patch0_patch_data_offset);
 }
 
 static LLVMValueRef get_num_tcs_out_vertices(struct si_shader_context *ctx)
@@ -415,14 +393,14 @@ static LLVMValueRef get_tcs_in_vertex_dw_stride(struct si_shader_context *ctx)
 
 	switch (ctx->type) {
 	case PIPE_SHADER_VERTEX:
-		stride = util_last_bit64(ctx->shader->selector->outputs_written);
-		return LLVMConstInt(ctx->i32, stride * 4, 0);
+		stride = ctx->shader->selector->lshs_vertex_stride / 4;
+		return LLVMConstInt(ctx->i32, stride, 0);
 
 	case PIPE_SHADER_TESS_CTRL:
 		if (ctx->screen->info.chip_class >= GFX9 &&
 		    ctx->shader->is_monolithic) {
-			stride = util_last_bit64(ctx->shader->key.part.tcs.ls->outputs_written);
-			return LLVMConstInt(ctx->i32, stride * 4, 0);
+			stride = ctx->shader->key.part.tcs.ls->lshs_vertex_stride / 4;
+			return LLVMConstInt(ctx->i32, stride, 0);
 		}
 		return si_unpack_param(ctx, ctx->param_vs_state_bits, 24, 8);
 
@@ -432,35 +410,6 @@ static LLVMValueRef get_tcs_in_vertex_dw_stride(struct si_shader_context *ctx)
 	}
 }
 
-static LLVMValueRef get_instance_index_for_fetch(
-	struct si_shader_context *ctx,
-	unsigned param_start_instance, LLVMValueRef divisor)
-{
-	LLVMValueRef result = ctx->abi.instance_id;
-
-	/* The division must be done before START_INSTANCE is added. */
-	if (divisor != ctx->i32_1)
-		result = LLVMBuildUDiv(ctx->ac.builder, result, divisor, "");
-
-	return LLVMBuildAdd(ctx->ac.builder, result,
-			    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->ac.builder;
-	LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->ac.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 LLVMValueRef unpack_sint16(struct si_shader_context *ctx,
 				 LLVMValueRef i32, unsigned index)
 {
@@ -552,17 +501,12 @@ void si_llvm_load_input_vs(
 		return;
 	}
 
-	unsigned chan;
-	unsigned fix_fetch;
-	unsigned num_fetches;
-	unsigned fetch_stride;
-	unsigned num_channels;
-
+	union si_vs_fix_fetch fix_fetch;
 	LLVMValueRef t_list_ptr;
 	LLVMValueRef t_offset;
 	LLVMValueRef t_list;
 	LLVMValueRef vertex_index;
-	LLVMValueRef input[3];
+	LLVMValueRef tmp;
 
 	/* Load the T list */
 	t_list_ptr = LLVMGetParam(ctx->main_fn, ctx->param_vertex_buffers);
@@ -575,66 +519,84 @@ void si_llvm_load_input_vs(
 				    ctx->param_vertex_index0 +
 				    input_index);
 
-	fix_fetch = ctx->shader->key.mono.vs_fix_fetch[input_index];
+	/* Use the open-coded implementation for all loads of doubles and
+	 * of dword-sized data that needs fixups. We need to insert conversion
+	 * code anyway, and the amd/common code does it for us.
+	 *
+	 * Note: On LLVM <= 8, we can only open-code formats with
+	 * channel size >= 4 bytes.
+	 */
+	bool opencode = ctx->shader->key.mono.vs_fetch_opencode & (1 << input_index);
+	fix_fetch.bits = ctx->shader->key.mono.vs_fix_fetch[input_index].bits;
+	if (opencode ||
+	    (fix_fetch.u.log_size == 3 && fix_fetch.u.format == AC_FETCH_FORMAT_FLOAT) ||
+	    (fix_fetch.u.log_size == 2)) {
+		tmp = ac_build_opencoded_load_format(
+				&ctx->ac, fix_fetch.u.log_size, fix_fetch.u.num_channels_m1 + 1,
+				fix_fetch.u.format, fix_fetch.u.reverse, !opencode,
+				t_list, vertex_index, ctx->ac.i32_0, ctx->ac.i32_0,
+				false, false, true);
+		for (unsigned i = 0; i < 4; ++i)
+			out[i] = LLVMBuildExtractElement(ctx->ac.builder, tmp, LLVMConstInt(ctx->i32, i, false), "");
+		return;
+	}
 
 	/* 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;
-		num_channels = 2;
-		break;
-	case SI_FIX_FETCH_RGBA_64_FLOAT:
-		num_fetches = 2; /* 2 4-dword loads */
-		fetch_stride = 16;
-		num_channels = 4;
-		break;
-	case SI_FIX_FETCH_RGB_8:
-	case SI_FIX_FETCH_RGB_8_INT:
-		num_fetches = 3;
-		fetch_stride = 1;
-		num_channels = 1;
-		break;
-	case SI_FIX_FETCH_RGB_16:
-	case SI_FIX_FETCH_RGB_16_INT:
-		num_fetches = 3;
-		fetch_stride = 2;
-		num_channels = 1;
-		break;
-	default:
+	unsigned required_channels = util_last_bit(info->input_usage_mask[input_index]);
+	LLVMValueRef fetches[4];
+	unsigned num_fetches;
+	unsigned fetch_stride;
+	unsigned channels_per_fetch;
+
+	if (fix_fetch.u.log_size <= 1 && fix_fetch.u.num_channels_m1 == 2) {
+		num_fetches = MIN2(required_channels, 3);
+		fetch_stride = 1 << fix_fetch.u.log_size;
+		channels_per_fetch = 1;
+	} else {
 		num_fetches = 1;
 		fetch_stride = 0;
-		num_channels = util_last_bit(info->input_usage_mask[input_index]);
+		channels_per_fetch = required_channels;
 	}
 
-	for (unsigned i = 0; i < num_fetches; i++) {
+	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,
-						       num_channels, false, true);
-		input[i] = ac_build_expand_to_vec4(&ctx->ac, input[i], num_channels);
+		fetches[i] = ac_build_buffer_load_format(&ctx->ac, t_list, vertex_index, voffset,
+							 channels_per_fetch, false, true);
 	}
 
-	/* Break up the vec4 into individual components */
-	for (chan = 0; chan < 4; chan++) {
-		LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, 0);
-		out[chan] = LLVMBuildExtractElement(ctx->ac.builder,
-						    input[0], llvm_chan, "");
+	if (num_fetches == 1 && channels_per_fetch > 1) {
+		LLVMValueRef fetch = fetches[0];
+		for (unsigned i = 0; i < channels_per_fetch; ++i) {
+			tmp = LLVMConstInt(ctx->i32, i, false);
+			fetches[i] = LLVMBuildExtractElement(
+				ctx->ac.builder, fetch, tmp, "");
+		}
+		num_fetches = channels_per_fetch;
+		channels_per_fetch = 1;
 	}
 
-	switch (fix_fetch) {
-	case SI_FIX_FETCH_A2_SNORM:
-	case SI_FIX_FETCH_A2_SSCALED:
-	case SI_FIX_FETCH_A2_SINT: {
-		/* The hardware returns an unsigned value; convert it to a
-		 * signed one.
+	for (unsigned i = num_fetches; i < 4; ++i)
+		fetches[i] = LLVMGetUndef(ctx->f32);
+
+	if (fix_fetch.u.log_size <= 1 && fix_fetch.u.num_channels_m1 == 2 &&
+	    required_channels == 4) {
+		if (fix_fetch.u.format == AC_FETCH_FORMAT_UINT || fix_fetch.u.format == AC_FETCH_FORMAT_SINT)
+			fetches[3] = ctx->ac.i32_1;
+		else
+			fetches[3] = ctx->ac.f32_1;
+	} else if (fix_fetch.u.log_size == 3 &&
+		   (fix_fetch.u.format == AC_FETCH_FORMAT_SNORM ||
+		    fix_fetch.u.format == AC_FETCH_FORMAT_SSCALED ||
+		    fix_fetch.u.format == AC_FETCH_FORMAT_SINT) &&
+		   required_channels == 4) {
+		/* For 2_10_10_10, the hardware returns an unsigned value;
+		 * convert it to a signed one.
 		 */
-		LLVMValueRef tmp = out[3];
+		LLVMValueRef tmp = fetches[3];
 		LLVMValueRef c30 = LLVMConstInt(ctx->i32, 30, 0);
 
 		/* First, recover the sign-extended signed integer value. */
-		if (fix_fetch == SI_FIX_FETCH_A2_SSCALED)
+		if (fix_fetch.u.format == AC_FETCH_FORMAT_SSCALED)
 			tmp = LLVMBuildFPToUI(ctx->ac.builder, tmp, ctx->i32, "");
 		else
 			tmp = ac_to_integer(&ctx->ac, tmp);
@@ -646,110 +608,26 @@ void si_llvm_load_input_vs(
 		 * exponent.
 		 */
 		tmp = LLVMBuildShl(ctx->ac.builder, tmp,
-				   fix_fetch == SI_FIX_FETCH_A2_SNORM ?
+				   fix_fetch.u.format == AC_FETCH_FORMAT_SNORM ?
 				   LLVMConstInt(ctx->i32, 7, 0) : c30, "");
 		tmp = LLVMBuildAShr(ctx->ac.builder, tmp, c30, "");
 
 		/* Convert back to the right type. */
-		if (fix_fetch == SI_FIX_FETCH_A2_SNORM) {
+		if (fix_fetch.u.format == AC_FETCH_FORMAT_SNORM) {
 			LLVMValueRef clamp;
 			LLVMValueRef neg_one = LLVMConstReal(ctx->f32, -1.0);
 			tmp = LLVMBuildSIToFP(ctx->ac.builder, tmp, ctx->f32, "");
 			clamp = LLVMBuildFCmp(ctx->ac.builder, LLVMRealULT, tmp, neg_one, "");
 			tmp = LLVMBuildSelect(ctx->ac.builder, clamp, neg_one, tmp, "");
-		} else if (fix_fetch == SI_FIX_FETCH_A2_SSCALED) {
+		} else if (fix_fetch.u.format == AC_FETCH_FORMAT_SSCALED) {
 			tmp = LLVMBuildSIToFP(ctx->ac.builder, tmp, ctx->f32, "");
 		}
 
-		out[3] = tmp;
-		break;
+		fetches[3] = tmp;
 	}
-	case SI_FIX_FETCH_RGBA_32_UNORM:
-	case SI_FIX_FETCH_RGBX_32_UNORM:
-		for (chan = 0; chan < 4; chan++) {
-			out[chan] = ac_to_integer(&ctx->ac, out[chan]);
-			out[chan] = LLVMBuildUIToFP(ctx->ac.builder,
-						    out[chan], ctx->f32, "");
-			out[chan] = LLVMBuildFMul(ctx->ac.builder, out[chan],
-						  LLVMConstReal(ctx->f32, 1.0 / UINT_MAX), "");
-		}
-		/* RGBX UINT returns 1 in alpha, which would be rounded to 0 by normalizing. */
-		if (fix_fetch == SI_FIX_FETCH_RGBX_32_UNORM)
-			out[3] = LLVMConstReal(ctx->f32, 1);
-		break;
-	case SI_FIX_FETCH_RGBA_32_SNORM:
-	case SI_FIX_FETCH_RGBX_32_SNORM:
-	case SI_FIX_FETCH_RGBA_32_FIXED:
-	case SI_FIX_FETCH_RGBX_32_FIXED: {
-		double scale;
-		if (fix_fetch >= SI_FIX_FETCH_RGBA_32_FIXED)
-			scale = 1.0 / 0x10000;
-		else
-			scale = 1.0 / INT_MAX;
-
-		for (chan = 0; chan < 4; chan++) {
-			out[chan] = ac_to_integer(&ctx->ac, out[chan]);
-			out[chan] = LLVMBuildSIToFP(ctx->ac.builder,
-						    out[chan], ctx->f32, "");
-			out[chan] = LLVMBuildFMul(ctx->ac.builder, out[chan],
-						  LLVMConstReal(ctx->f32, scale), "");
-		}
-		/* RGBX SINT returns 1 in alpha, which would be rounded to 0 by normalizing. */
-		if (fix_fetch == SI_FIX_FETCH_RGBX_32_SNORM ||
-		    fix_fetch == SI_FIX_FETCH_RGBX_32_FIXED)
-			out[3] = LLVMConstReal(ctx->f32, 1);
-		break;
-	}
-	case SI_FIX_FETCH_RGBA_32_USCALED:
-		for (chan = 0; chan < 4; chan++) {
-			out[chan] = ac_to_integer(&ctx->ac, out[chan]);
-			out[chan] = LLVMBuildUIToFP(ctx->ac.builder,
-						    out[chan], ctx->f32, "");
-		}
-		break;
-	case SI_FIX_FETCH_RGBA_32_SSCALED:
-		for (chan = 0; chan < 4; chan++) {
-			out[chan] = ac_to_integer(&ctx->ac, out[chan]);
-			out[chan] = LLVMBuildSIToFP(ctx->ac.builder,
-						    out[chan], ctx->f32, "");
-		}
-		break;
-	case SI_FIX_FETCH_RG_64_FLOAT:
-		for (chan = 0; chan < 2; chan++)
-			out[chan] = extract_double_to_float(ctx, input[0], chan);
-
-		out[2] = LLVMConstReal(ctx->f32, 0);
-		out[3] = LLVMConstReal(ctx->f32, 1);
-		break;
-	case SI_FIX_FETCH_RGB_64_FLOAT:
-		for (chan = 0; chan < 3; chan++)
-			out[chan] = extract_double_to_float(ctx, input[chan], 0);
 
-		out[3] = LLVMConstReal(ctx->f32, 1);
-		break;
-	case SI_FIX_FETCH_RGBA_64_FLOAT:
-		for (chan = 0; chan < 4; chan++) {
-			out[chan] = extract_double_to_float(ctx, input[chan / 2],
-							    chan % 2);
-		}
-		break;
-	case SI_FIX_FETCH_RGB_8:
-	case SI_FIX_FETCH_RGB_8_INT:
-	case SI_FIX_FETCH_RGB_16:
-	case SI_FIX_FETCH_RGB_16_INT:
-		for (chan = 0; chan < 3; chan++) {
-			out[chan] = LLVMBuildExtractElement(ctx->ac.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] = ac_to_float(&ctx->ac, ctx->i32_1);
-		}
-		break;
-	}
+	for (unsigned i = 0; i < 4; ++i)
+		out[i] = ac_to_float(&ctx->ac, fetches[i]);
 }
 
 static void declare_input_vs(
@@ -761,8 +639,8 @@ static void declare_input_vs(
 	si_llvm_load_input_vs(ctx, input_index, out);
 }
 
-static LLVMValueRef get_primitive_id(struct si_shader_context *ctx,
-				     unsigned swizzle)
+LLVMValueRef si_get_primitive_id(struct si_shader_context *ctx,
+				 unsigned swizzle)
 {
 	if (swizzle > 0)
 		return ctx->i32_0;
@@ -813,12 +691,8 @@ LLVMValueRef si_get_indirect_index(struct si_shader_context *ctx,
 		result = ac_to_integer(&ctx->ac, result);
 	}
 
-	if (addr_mul != 1)
-		result = LLVMBuildMul(ctx->ac.builder, result,
-				      LLVMConstInt(ctx->i32, addr_mul, 0), "");
-	result = LLVMBuildAdd(ctx->ac.builder, result,
-			      LLVMConstInt(ctx->i32, rel_index, 0), "");
-	return result;
+	return ac_build_imad(&ctx->ac, result, LLVMConstInt(ctx->i32, addr_mul, 0),
+			     LLVMConstInt(ctx->i32, rel_index, 0));
 }
 
 /**
@@ -845,22 +719,20 @@ static LLVMValueRef get_dw_address_from_generic_indices(struct si_shader_context
 							bool is_patch)
 {
 	if (vertex_dw_stride) {
-		base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
-					 LLVMBuildMul(ctx->ac.builder, vertex_index,
-						      vertex_dw_stride, ""), "");
+		base_addr = ac_build_imad(&ctx->ac, vertex_index,
+					  vertex_dw_stride, base_addr);
 	}
 
 	if (param_index) {
-		base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
-					 LLVMBuildMul(ctx->ac.builder, param_index,
-						      LLVMConstInt(ctx->i32, 4, 0), ""), "");
+		base_addr = ac_build_imad(&ctx->ac, param_index,
+					  LLVMConstInt(ctx->i32, 4, 0), base_addr);
 	}
 
 	int param = is_patch ?
 		si_shader_io_get_unique_index_patch(name[input_index],
 						    index[input_index]) :
 		si_shader_io_get_unique_index(name[input_index],
-					      index[input_index]);
+					      index[input_index], false);
 
 	/* Add the base address of the element. */
 	return LLVMBuildAdd(ctx->ac.builder, base_addr,
@@ -973,22 +845,15 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx,
 
 	constant16 = LLVMConstInt(ctx->i32, 16, 0);
 	if (vertex_index) {
-		base_addr = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
-		                         vertices_per_patch, "");
-
-		base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
-		                         vertex_index, "");
-
+		base_addr = ac_build_imad(&ctx->ac, rel_patch_id,
+					  vertices_per_patch, vertex_index);
 		param_stride = total_vertices;
 	} else {
 		base_addr = rel_patch_id;
 		param_stride = num_patches;
 	}
 
-	base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
-	                         LLVMBuildMul(ctx->ac.builder, param_index,
-	                                      param_stride, ""), "");
-
+	base_addr = ac_build_imad(&ctx->ac, param_index, param_stride, base_addr);
 	base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, "");
 
 	if (!vertex_index) {
@@ -1015,7 +880,7 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_generic_indices(
 
 	param_index_base = is_patch ?
 		si_shader_io_get_unique_index_patch(name[param_base], index[param_base]) :
-		si_shader_io_get_unique_index(name[param_base], index[param_base]);
+		si_shader_io_get_unique_index(name[param_base], index[param_base], false);
 
 	if (param_index) {
 		param_index = LLVMBuildAdd(ctx->ac.builder, param_index,
@@ -1119,13 +984,13 @@ static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
 }
 
 /**
- * Load from LDS.
+ * Load from LSHS LDS storage.
  *
  * \param type		output value type
  * \param swizzle	offset (typically 0..3); it can be ~0, which loads a vec4
  * \param dw_addr	address in dwords
  */
-static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
+static LLVMValueRef lshs_lds_load(struct lp_build_tgsi_context *bld_base,
 			     LLVMTypeRef type, unsigned swizzle,
 			     LLVMValueRef dw_addr)
 {
@@ -1136,9 +1001,9 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
 		LLVMValueRef values[TGSI_NUM_CHANNELS];
 
 		for (unsigned chan = 0; chan < TGSI_NUM_CHANNELS; chan++)
-			values[chan] = lds_load(bld_base, type, chan, dw_addr);
+			values[chan] = lshs_lds_load(bld_base, type, chan, dw_addr);
 
-		return lp_build_gather_values(&ctx->gallivm, values,
+		return ac_build_gather_values(&ctx->ac, values,
 					      TGSI_NUM_CHANNELS);
 	}
 
@@ -1146,13 +1011,13 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
 	if (llvm_type_is_64bit(ctx, type)) {
 		LLVMValueRef lo, hi;
 
-		lo = lds_load(bld_base, ctx->i32, swizzle, dw_addr);
-		hi = lds_load(bld_base, ctx->i32, swizzle + 1, dw_addr);
+		lo = lshs_lds_load(bld_base, ctx->i32, swizzle, dw_addr);
+		hi = lshs_lds_load(bld_base, ctx->i32, swizzle + 1, dw_addr);
 		return si_llvm_emit_fetch_64bit(bld_base, type, lo, hi);
 	}
 
-	dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
-			    LLVMConstInt(ctx->i32, swizzle, 0));
+	dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
+			       LLVMConstInt(ctx->i32, swizzle, 0), "");
 
 	value = ac_lds_load(&ctx->ac, dw_addr);
 
@@ -1160,18 +1025,18 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
 }
 
 /**
- * Store to LDS.
+ * Store to LSHS LDS storage.
  *
  * \param swizzle	offset (typically 0..3)
  * \param dw_addr	address in dwords
  * \param value		value to store
  */
-static void lds_store(struct si_shader_context *ctx,
+static void lshs_lds_store(struct si_shader_context *ctx,
 		      unsigned dw_offset_imm, LLVMValueRef dw_addr,
 		      LLVMValueRef value)
 {
-	dw_addr = lp_build_add(&ctx->bld_base.uint_bld, dw_addr,
-			    LLVMConstInt(ctx->i32, dw_offset_imm, 0));
+	dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
+			       LLVMConstInt(ctx->i32, dw_offset_imm, 0), "");
 
 	ac_lds_store(&ctx->ac, dw_addr, value);
 }
@@ -1221,16 +1086,16 @@ static LLVMValueRef get_tess_ring_descriptor(struct si_shader_context *ctx,
 static LLVMValueRef fetch_input_tcs(
 	struct lp_build_tgsi_context *bld_base,
 	const struct tgsi_full_src_register *reg,
-	enum tgsi_opcode_type type, unsigned swizzle)
+	enum tgsi_opcode_type type, unsigned swizzle_in)
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
 	LLVMValueRef dw_addr, stride;
-
+	unsigned swizzle = swizzle_in & 0xffff;
 	stride = get_tcs_in_vertex_dw_stride(ctx);
 	dw_addr = get_tcs_in_current_patch_offset(ctx);
 	dw_addr = get_dw_address(ctx, NULL, reg, stride, dw_addr);
 
-	return lds_load(bld_base, tgsi2llvmtype(bld_base, type), swizzle, dw_addr);
+	return lshs_lds_load(bld_base, tgsi2llvmtype(bld_base, type), swizzle, dw_addr);
 }
 
 static LLVMValueRef si_nir_load_tcs_varyings(struct ac_shader_abi *abi,
@@ -1297,7 +1162,7 @@ static LLVMValueRef si_nir_load_tcs_varyings(struct ac_shader_abi *abi,
 			offset *= 2;
 
 		offset += component;
-		value[i + component] = lds_load(bld_base, type, offset, dw_addr);
+		value[i + component] = lshs_lds_load(bld_base, type, offset, dw_addr);
 	}
 
 	return ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
@@ -1306,10 +1171,11 @@ static LLVMValueRef si_nir_load_tcs_varyings(struct ac_shader_abi *abi,
 static LLVMValueRef fetch_output_tcs(
 		struct lp_build_tgsi_context *bld_base,
 		const struct tgsi_full_src_register *reg,
-		enum tgsi_opcode_type type, unsigned swizzle)
+		enum tgsi_opcode_type type, unsigned swizzle_in)
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
 	LLVMValueRef dw_addr, stride;
+	unsigned swizzle = (swizzle_in & 0xffff);
 
 	if (reg->Register.Dimension) {
 		stride = get_tcs_out_vertex_dw_stride(ctx);
@@ -1320,16 +1186,17 @@ static LLVMValueRef fetch_output_tcs(
 		dw_addr = get_dw_address(ctx, NULL, reg, NULL, dw_addr);
 	}
 
-	return lds_load(bld_base, tgsi2llvmtype(bld_base, type), swizzle, dw_addr);
+	return lshs_lds_load(bld_base, tgsi2llvmtype(bld_base, type), swizzle, dw_addr);
 }
 
 static LLVMValueRef fetch_input_tes(
 	struct lp_build_tgsi_context *bld_base,
 	const struct tgsi_full_src_register *reg,
-	enum tgsi_opcode_type type, unsigned swizzle)
+	enum tgsi_opcode_type type, unsigned swizzle_in)
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
 	LLVMValueRef base, addr;
+	unsigned swizzle = (swizzle_in & 0xffff);
 
 	base = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
 	addr = get_tcs_tes_buffer_address_from_reg(ctx, NULL, reg);
@@ -1381,8 +1248,20 @@ LLVMValueRef si_nir_load_input_tes(struct ac_shader_abi *abi,
 	LLVMValueRef value[4];
 	for (unsigned i = 0; i < num_components; i++) {
 		unsigned offset = i;
-		if (llvm_type_is_64bit(ctx, type))
+		if (llvm_type_is_64bit(ctx, type)) {
 			offset *= 2;
+			if (offset == 4) {
+                                addr = get_tcs_tes_buffer_address_from_generic_indices(ctx,
+                                                                                       vertex_index,
+                                                                                       param_index,
+                                                                                       driver_location + 1,
+                                                                                       info->input_semantic_name,
+                                                                                       info->input_semantic_index,
+                                                                                       is_patch);
+			}
+
+                        offset = offset % 4;
+		}
 
 		offset += component;
 		value[i + component] = buffer_load(&ctx->bld_base, type, offset,
@@ -1457,7 +1336,7 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
 
 		/* Skip LDS stores if there is no LDS read of this output. */
 		if (!skip_lds_store)
-			lds_store(ctx, chan_index, dw_addr, value);
+			lshs_lds_store(ctx, chan_index, dw_addr, value);
 
 		value = ac_to_integer(&ctx->ac, value);
 		values[chan_index] = value;
@@ -1465,7 +1344,7 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
 		if (reg->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);
+						    4 * chan_index, 1, 0, false);
 		}
 
 		/* Write tess factors into VGPRs for the epilog. */
@@ -1482,10 +1361,10 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
 	}
 
 	if (reg->Register.WriteMask == 0xF && !is_tess_factor) {
-		LLVMValueRef value = lp_build_gather_values(&ctx->gallivm,
+		LLVMValueRef value = ac_build_gather_values(&ctx->ac,
 		                                            values, 4);
 		ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buf_addr,
-					    base, 0, 1, 0, true, false);
+					    base, 0, 1, 0, false);
 	}
 }
 
@@ -1504,7 +1383,7 @@ static void si_nir_store_output_tcs(struct ac_shader_abi *abi,
 	unsigned driver_location = var->data.driver_location;
 	LLVMValueRef dw_addr, stride;
 	LLVMValueRef buffer, base, addr;
-	LLVMValueRef values[4];
+	LLVMValueRef values[8];
 	bool skip_lds_store;
 	bool is_tess_factor = false, is_tess_inner = false;
 
@@ -1566,14 +1445,25 @@ static void si_nir_store_output_tcs(struct ac_shader_abi *abi,
 							       info->output_semantic_index,
 							       is_patch);
 
-	for (unsigned chan = 0; chan < 4; chan++) {
+	for (unsigned chan = 0; chan < 8; chan++) {
 		if (!(writemask & (1 << chan)))
 			continue;
 		LLVMValueRef value = ac_llvm_extract_elem(&ctx->ac, src, chan - component);
 
+		unsigned buffer_store_offset = chan % 4;
+		if (chan == 4) {
+                        addr = get_tcs_tes_buffer_address_from_generic_indices(ctx,
+                                                                               vertex_index,
+                                                                               param_index,
+                                                                               driver_location + 1,
+                                                                               info->output_semantic_name,
+                                                                               info->output_semantic_index,
+                                                                               is_patch);
+		}
+
 		/* Skip LDS stores if there is no LDS read of this output. */
 		if (!skip_lds_store)
-			lds_store(ctx, chan, dw_addr, value);
+			lshs_lds_store(ctx, chan, dw_addr, value);
 
 		value = ac_to_integer(&ctx->ac, value);
 		values[chan] = value;
@@ -1581,7 +1471,8 @@ static void si_nir_store_output_tcs(struct ac_shader_abi *abi,
 		if (writemask != 0xF && !is_tess_factor) {
 			ac_build_buffer_store_dword(&ctx->ac, buffer, value, 1,
 						    addr, base,
-						    4 * chan, 1, 0, true, false);
+						    4 * buffer_store_offset,
+                                                    1, 0, false);
 		}
 
 		/* Write tess factors into VGPRs for the epilog. */
@@ -1598,10 +1489,10 @@ static void si_nir_store_output_tcs(struct ac_shader_abi *abi,
 	}
 
 	if (writemask == 0xF && !is_tess_factor) {
-		LLVMValueRef value = lp_build_gather_values(&ctx->gallivm,
+		LLVMValueRef value = ac_build_gather_values(&ctx->ac,
 		                                            values, 4);
 		ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, addr,
-					    base, 0, 1, 0, true, false);
+					    base, 0, 1, 0, false);
 	}
 }
 
@@ -1614,7 +1505,6 @@ LLVMValueRef si_llvm_load_input_gs(struct ac_shader_abi *abi,
 	struct si_shader_context *ctx = si_shader_context_from_abi(abi);
 	struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
 	struct si_shader *shader = ctx->shader;
-	struct lp_build_context *uint =	&ctx->bld_base.uint_bld;
 	LLVMValueRef vtx_offset, soffset;
 	struct tgsi_shader_info *info = &shader->selector->info;
 	unsigned semantic_name = info->input_semantic_name[input_index];
@@ -1622,7 +1512,7 @@ LLVMValueRef si_llvm_load_input_gs(struct ac_shader_abi *abi,
 	unsigned param;
 	LLVMValueRef value;
 
-	param = si_shader_io_get_unique_index(semantic_name, semantic_index);
+	param = si_shader_io_get_unique_index(semantic_name, semantic_index, false);
 
 	/* GFX9 has the ESGS ring in LDS. */
 	if (ctx->screen->info.chip_class >= GFX9) {
@@ -1646,9 +1536,22 @@ LLVMValueRef si_llvm_load_input_gs(struct ac_shader_abi *abi,
 			return NULL;
 		}
 
+		unsigned offset = param * 4 + swizzle;
 		vtx_offset = LLVMBuildAdd(ctx->ac.builder, vtx_offset,
-					  LLVMConstInt(ctx->i32, param * 4, 0), "");
-		return lds_load(bld_base, type, swizzle, vtx_offset);
+					  LLVMConstInt(ctx->i32, offset, false), "");
+
+		LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->esgs_ring, vtx_offset);
+		LLVMValueRef value = LLVMBuildLoad(ctx->ac.builder, ptr, "");
+		if (llvm_type_is_64bit(ctx, type)) {
+			ptr = LLVMBuildGEP(ctx->ac.builder, ptr,
+					   &ctx->ac.i32_1, 1, "");
+			LLVMValueRef values[2] = {
+				value,
+				LLVMBuildLoad(ctx->ac.builder, ptr, "")
+			};
+			value = ac_build_gather_values(&ctx->ac, values, 2);
+		}
+		return LLVMBuildBitCast(ctx->ac.builder, value, type, "");
 	}
 
 	/* GFX6: input load from the ESGS ring in memory. */
@@ -1659,14 +1562,15 @@ LLVMValueRef si_llvm_load_input_gs(struct ac_shader_abi *abi,
 			values[chan] = si_llvm_load_input_gs(abi, input_index, vtx_offset_param,
 							     type, chan);
 		}
-		return lp_build_gather_values(&ctx->gallivm, values,
+		return ac_build_gather_values(&ctx->ac, values,
 					      TGSI_NUM_CHANNELS);
 	}
 
 	/* Get the vertex offset parameter on GFX6. */
 	LLVMValueRef gs_vtx_offset = ctx->gs_vtx_offset[vtx_offset_param];
 
-	vtx_offset = lp_build_mul_imm(uint, gs_vtx_offset, 4);
+	vtx_offset = LLVMBuildMul(ctx->ac.builder, gs_vtx_offset,
+				  LLVMConstInt(ctx->i32, 4, 0), "");
 
 	soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle) * 256, 0);
 
@@ -1713,14 +1617,15 @@ static LLVMValueRef fetch_input_gs(
 	struct lp_build_tgsi_context *bld_base,
 	const struct tgsi_full_src_register *reg,
 	enum tgsi_opcode_type type,
-	unsigned swizzle)
+	unsigned swizzle_in)
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
 	struct tgsi_shader_info *info = &ctx->shader->selector->info;
+	unsigned swizzle = swizzle_in & 0xffff;
 
 	unsigned semantic_name = info->input_semantic_name[reg->Register.Index];
 	if (swizzle != ~0 && semantic_name == TGSI_SEMANTIC_PRIMID)
-		return get_primitive_id(ctx, swizzle);
+		return si_get_primitive_id(ctx, swizzle);
 
 	if (!reg->Register.Dimension)
 		return NULL;
@@ -1881,7 +1786,6 @@ void si_llvm_load_input_fs(
 	unsigned input_index,
 	LLVMValueRef out[4])
 {
-	struct lp_build_context *base = &ctx->bld_base.base;
 	struct si_shader *shader = ctx->shader;
 	struct tgsi_shader_info *info = &shader->selector->info;
 	LLVMValueRef main_fn = ctx->main_fn;
@@ -1898,11 +1802,12 @@ void si_llvm_load_input_fs(
 		unsigned mask = colors_read >> (semantic_index * 4);
 		unsigned offset = SI_PARAM_POS_FIXED_PT + 1 +
 				  (semantic_index ? util_bitcount(colors_read & 0xf) : 0);
+		LLVMValueRef undef = LLVMGetUndef(ctx->f32);
 
-		out[0] = mask & 0x1 ? LLVMGetParam(main_fn, offset++) : base->undef;
-		out[1] = mask & 0x2 ? LLVMGetParam(main_fn, offset++) : base->undef;
-		out[2] = mask & 0x4 ? LLVMGetParam(main_fn, offset++) : base->undef;
-		out[3] = mask & 0x8 ? LLVMGetParam(main_fn, offset++) : base->undef;
+		out[0] = mask & 0x1 ? LLVMGetParam(main_fn, offset++) : undef;
+		out[1] = mask & 0x2 ? LLVMGetParam(main_fn, offset++) : undef;
+		out[2] = mask & 0x4 ? LLVMGetParam(main_fn, offset++) : undef;
+		out[3] = mask & 0x8 ? LLVMGetParam(main_fn, offset++) : undef;
 		return;
 	}
 
@@ -1973,7 +1878,7 @@ static LLVMValueRef get_block_size(struct ac_shader_abi *abi)
 		for (i = 0; i < 3; ++i)
 			values[i] = LLVMConstInt(ctx->i32, sizes[i], 0);
 
-		result = lp_build_gather_values(&ctx->gallivm, values, 3);
+		result = ac_build_gather_values(&ctx->ac, values, 3);
 	} else {
 		result = LLVMGetParam(ctx->main_fn, ctx->param_block_size);
 	}
@@ -1995,13 +1900,12 @@ static LLVMValueRef buffer_load_const(struct si_shader_context *ctx,
 static LLVMValueRef load_sample_position(struct ac_shader_abi *abi, LLVMValueRef sample_id)
 {
 	struct si_shader_context *ctx = si_shader_context_from_abi(abi);
-	struct lp_build_context *uint_bld = &ctx->bld_base.uint_bld;
 	LLVMValueRef desc = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
 	LLVMValueRef buf_index = LLVMConstInt(ctx->i32, SI_PS_CONST_SAMPLE_POSITIONS, 0);
 	LLVMValueRef resource = ac_build_load_to_sgpr(&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 offset0 = LLVMBuildMul(ctx->ac.builder, sample_id, LLVMConstInt(ctx->i32, 8, 0), "");
 	LLVMValueRef offset1 = LLVMBuildAdd(ctx->ac.builder, offset0, LLVMConstInt(ctx->i32, 4, 0), "");
 
 	LLVMValueRef pos[4] = {
@@ -2011,19 +1915,18 @@ static LLVMValueRef load_sample_position(struct ac_shader_abi *abi, LLVMValueRef
 		LLVMConstReal(ctx->f32, 0)
 	};
 
-	return lp_build_gather_values(&ctx->gallivm, pos, 4);
+	return ac_build_gather_values(&ctx->ac, pos, 4);
 }
 
 static LLVMValueRef load_sample_mask_in(struct ac_shader_abi *abi)
 {
-	return abi->sample_coverage;
+	struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+	return ac_to_integer(&ctx->ac, abi->sample_coverage);
 }
 
 static LLVMValueRef si_load_tess_coord(struct ac_shader_abi *abi)
 {
 	struct si_shader_context *ctx = si_shader_context_from_abi(abi);
-	struct lp_build_context *bld = &ctx->bld_base.base;
-
 	LLVMValueRef coord[4] = {
 		LLVMGetParam(ctx->main_fn, ctx->param_tes_u),
 		LLVMGetParam(ctx->main_fn, ctx->param_tes_v),
@@ -2033,11 +1936,12 @@ static LLVMValueRef si_load_tess_coord(struct ac_shader_abi *abi)
 
 	/* For triangles, the vector should be (u, v, 1-u-v). */
 	if (ctx->shader->selector->info.properties[TGSI_PROPERTY_TES_PRIM_MODE] ==
-	    PIPE_PRIM_TRIANGLES)
-		coord[2] = lp_build_sub(bld, ctx->ac.f32_1,
-					lp_build_add(bld, coord[0], coord[1]));
-
-	return lp_build_gather_values(&ctx->gallivm, coord, 4);
+	    PIPE_PRIM_TRIANGLES) {
+		coord[2] = LLVMBuildFSub(ctx->ac.builder, ctx->ac.f32_1,
+					 LLVMBuildFAdd(ctx->ac.builder,
+						       coord[0], coord[1], ""), "");
+	}
+	return ac_build_gather_values(&ctx->ac, coord, 4);
 }
 
 static LLVMValueRef load_tess_level(struct si_shader_context *ctx,
@@ -2140,11 +2044,10 @@ void si_load_system_value(struct si_shader_context *ctx,
 			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)),
+			ac_build_fdiv(&ctx->ac, ctx->ac.f32_1,
+				      LLVMGetParam(ctx->main_fn, SI_PARAM_POS_W_FLOAT)),
 		};
-		value = lp_build_gather_values(&ctx->gallivm, pos, 4);
+		value = ac_build_gather_values(&ctx->ac, pos, 4);
 		break;
 	}
 
@@ -2163,11 +2066,9 @@ void si_load_system_value(struct si_shader_context *ctx,
 			LLVMConstReal(ctx->f32, 0),
 			LLVMConstReal(ctx->f32, 0)
 		};
-		pos[0] = lp_build_emit_llvm_unary(&ctx->bld_base,
-						  TGSI_OPCODE_FRC, pos[0]);
-		pos[1] = lp_build_emit_llvm_unary(&ctx->bld_base,
-						  TGSI_OPCODE_FRC, pos[1]);
-		value = lp_build_gather_values(&ctx->gallivm, pos, 4);
+		pos[0] = ac_build_fract(&ctx->ac, pos[0], 32);
+		pos[1] = ac_build_fract(&ctx->ac, pos[1], 32);
+		value = ac_build_gather_values(&ctx->ac, pos, 4);
 		break;
 	}
 
@@ -2205,12 +2106,12 @@ void si_load_system_value(struct si_shader_context *ctx,
 		for (i = 0; i < 4; i++)
 			val[i] = buffer_load_const(ctx, buf,
 						   LLVMConstInt(ctx->i32, (offset + i) * 4, 0));
-		value = lp_build_gather_values(&ctx->gallivm, val, 4);
+		value = ac_build_gather_values(&ctx->ac, val, 4);
 		break;
 	}
 
 	case TGSI_SEMANTIC_PRIMID:
-		value = get_primitive_id(ctx, 0);
+		value = si_get_primitive_id(ctx, 0);
 		break;
 
 	case TGSI_SEMANTIC_GRID_SIZE:
@@ -2231,7 +2132,7 @@ void si_load_system_value(struct si_shader_context *ctx,
 				values[i] = ctx->abi.workgroup_ids[i];
 			}
 		}
-		value = lp_build_gather_values(&ctx->gallivm, values, 3);
+		value = ac_build_gather_values(&ctx->ac, values, 3);
 		break;
 	}
 
@@ -2240,12 +2141,7 @@ void si_load_system_value(struct si_shader_context *ctx,
 		break;
 
 	case TGSI_SEMANTIC_HELPER_INVOCATION:
-		value = lp_build_intrinsic(ctx->ac.builder,
-					   "llvm.amdgcn.ps.live",
-					   ctx->i1, NULL, 0,
-					   LP_FUNC_ATTR_READNONE);
-		value = LLVMBuildNot(ctx->ac.builder, value, "");
-		value = LLVMBuildSExt(ctx->ac.builder, value, ctx->i32, "");
+		value = ac_build_load_helper_invocation(&ctx->ac);
 		break;
 
 	case TGSI_SEMANTIC_SUBGROUP_SIZE:
@@ -2288,6 +2184,10 @@ void si_load_system_value(struct si_shader_context *ctx,
 		break;
 	}
 
+	case TGSI_SEMANTIC_CS_USER_DATA:
+		value = LLVMGetParam(ctx->main_fn, ctx->param_cs_user_data);
+		break;
+
 	default:
 		assert(!"unknown system value");
 		return;
@@ -2299,17 +2199,18 @@ void si_load_system_value(struct si_shader_context *ctx,
 void si_declare_compute_memory(struct si_shader_context *ctx)
 {
 	struct si_shader_selector *sel = ctx->shader->selector;
+	unsigned lds_size = sel->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE];
 
-	LLVMTypeRef i8p = LLVMPointerType(ctx->i8, AC_LOCAL_ADDR_SPACE);
+	LLVMTypeRef i8p = LLVMPointerType(ctx->i8, AC_ADDR_SPACE_LDS);
 	LLVMValueRef var;
 
 	assert(!ctx->ac.lds);
 
 	var = LLVMAddGlobalInAddressSpace(ctx->ac.module,
-	                                  LLVMArrayType(ctx->i8, sel->local_size),
+	                                  LLVMArrayType(ctx->i8, lds_size),
 	                                  "compute_lds",
-	                                  AC_LOCAL_ADDR_SPACE);
-	LLVMSetAlignment(var, 4);
+	                                  AC_ADDR_SPACE_LDS);
+	LLVMSetAlignment(var, 64 * 1024);
 
 	ctx->ac.lds = LLVMBuildBitCast(ctx->ac.builder, var, i8p, "");
 }
@@ -2337,18 +2238,9 @@ static LLVMValueRef load_const_buffer_desc_fast_path(struct si_shader_context *c
 	ptr = LLVMBuildPtrToInt(ctx->ac.builder, ptr, ctx->ac.intptr, "");
 
 	LLVMValueRef desc0, desc1;
-	if (HAVE_32BIT_POINTERS) {
-		desc0 = ptr;
-		desc1 = LLVMConstInt(ctx->i32,
-				     S_008F04_BASE_ADDRESS_HI(ctx->screen->info.address32_hi), 0);
-	} else {
-		ptr = LLVMBuildBitCast(ctx->ac.builder, ptr, ctx->v2i32, "");
-		desc0 = LLVMBuildExtractElement(ctx->ac.builder, ptr, ctx->i32_0, "");
-		desc1 = LLVMBuildExtractElement(ctx->ac.builder, ptr, ctx->i32_1, "");
-		/* Mask out all bits except BASE_ADDRESS_HI. */
-		desc1 = LLVMBuildAnd(ctx->ac.builder, desc1,
-				     LLVMConstInt(ctx->i32, ~C_008F04_BASE_ADDRESS_HI, 0), "");
-	}
+	desc0 = ptr;
+	desc1 = LLVMConstInt(ctx->i32,
+			     S_008F04_BASE_ADDRESS_HI(ctx->screen->info.address32_hi), 0);
 
 	LLVMValueRef desc_elems[] = {
 		desc0,
@@ -2413,22 +2305,23 @@ static LLVMValueRef fetch_constant(
 	struct lp_build_tgsi_context *bld_base,
 	const struct tgsi_full_src_register *reg,
 	enum tgsi_opcode_type type,
-	unsigned swizzle)
+	unsigned swizzle_in)
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
 	struct si_shader_selector *sel = ctx->shader->selector;
 	const struct tgsi_ind_register *ireg = &reg->Indirect;
 	unsigned buf, idx;
+	unsigned swizzle = swizzle_in & 0xffff;
 
 	LLVMValueRef addr, bufp;
 
-	if (swizzle == LP_CHAN_ALL) {
+	if (swizzle_in == LP_CHAN_ALL) {
 		unsigned chan;
 		LLVMValueRef values[4];
 		for (chan = 0; chan < TGSI_NUM_CHANNELS; ++chan)
 			values[chan] = fetch_constant(bld_base, reg, type, chan);
 
-		return lp_build_gather_values(&ctx->gallivm, values, 4);
+		return ac_build_gather_values(&ctx->ac, values, 4);
 	}
 
 	/* Split 64-bit loads. */
@@ -2436,7 +2329,7 @@ static LLVMValueRef fetch_constant(
 		LLVMValueRef lo, hi;
 
 		lo = fetch_constant(bld_base, reg, TGSI_TYPE_UNSIGNED, swizzle);
-		hi = fetch_constant(bld_base, reg, TGSI_TYPE_UNSIGNED, swizzle + 1);
+		hi = fetch_constant(bld_base, reg, TGSI_TYPE_UNSIGNED, (swizzle_in >> 16));
 		return si_llvm_emit_fetch_64bit(bld_base, tgsi2llvmtype(bld_base, type),
 						lo, hi);
 	}
@@ -2451,30 +2344,6 @@ static LLVMValueRef fetch_constant(
 	/* Fast path when user data SGPRs point to constant buffer 0 directly. */
 	if (sel->info.const_buffers_declared == 1 &&
 	    sel->info.shader_buffers_declared == 0) {
-
-		/* This enables use of s_load_dword and flat_load_dword for const buffer 0
-		 * loads, and up to x4 load opcode merging. However, it leads to horrible
-		 * code reducing SIMD wave occupancy from 8 to 2 in many cases.
-		 *
-		 * Using s_buffer_load_dword (x1) seems to be the best option right now.
-		 *
-		 * LLVM 5.0 on SI doesn't insert a required s_nop between SALU setting
-		 * a descriptor and s_buffer_load_dword using it, so we can't expand
-		 * the pointer into a full descriptor like below. We have to use
-		 * s_load_dword instead. The only case when LLVM 5.0 would select
-		 * s_buffer_load_dword (that we have to prevent) is when we use use
-		 * a literal offset where we don't need bounds checking.
-		 */
-		if (ctx->screen->info.chip_class == SI && HAVE_LLVM < 0x0600 &&
-		    !reg->Register.Indirect) {
-			LLVMValueRef ptr =
-				LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers);
-
-			addr = LLVMBuildLShr(ctx->ac.builder, addr, LLVMConstInt(ctx->i32, 2, 0), "");
-			LLVMValueRef result = ac_build_load_invariant(&ctx->ac, ptr, addr);
-			return bitcast(bld_base, type, result);
-		}
-
 		LLVMValueRef desc = load_const_buffer_desc_fast_path(ctx);
 		LLVMValueRef result = buffer_load_const(ctx, desc, addr);
 		return bitcast(bld_base, type, result);
@@ -2645,7 +2514,7 @@ static void si_alpha_test(struct lp_build_tgsi_context *bld_base,
 			LLVMBuildFCmp(ctx->ac.builder, cond, alpha, alpha_ref, "");
 		ac_build_kill_if_false(&ctx->ac, alpha_pass);
 	} else {
-		ac_build_kill_if_false(&ctx->ac, LLVMConstInt(ctx->i1, 0, 0));
+		ac_build_kill_if_false(&ctx->ac, ctx->i1false);
 	}
 }
 
@@ -2661,9 +2530,9 @@ static LLVMValueRef si_scale_alpha_by_sample_mask(struct lp_build_tgsi_context *
 				samplemask_param);
 	coverage = ac_to_integer(&ctx->ac, coverage);
 
-	coverage = lp_build_intrinsic(ctx->ac.builder, "llvm.ctpop.i32",
+	coverage = ac_build_intrinsic(&ctx->ac, "llvm.ctpop.i32",
 				   ctx->i32,
-				   &coverage, 1, LP_FUNC_ATTR_READNONE);
+				   &coverage, 1, AC_FUNC_ATTR_READNONE);
 
 	coverage = LLVMBuildUIToFP(ctx->ac.builder, coverage,
 				   ctx->f32, "");
@@ -2703,10 +2572,8 @@ static void si_llvm_emit_clipvertex(struct si_shader_context *ctx,
 								const_chan) * 4, 0);
 				base_elt = buffer_load_const(ctx, const_resource,
 							     addr);
-				args->out[chan] =
-					lp_build_add(&ctx->bld_base.base, args->out[chan],
-						     lp_build_mul(&ctx->bld_base.base, base_elt,
-								  out_elts[const_chan]));
+				args->out[chan] = ac_build_fmad(&ctx->ac, base_elt,
+								out_elts[const_chan], args->out[chan]);
 			}
 		}
 
@@ -2739,11 +2606,11 @@ static void si_dump_streamout(struct pipe_stream_output_info *so)
 	}
 }
 
-static void emit_streamout_output(struct si_shader_context *ctx,
-				  LLVMValueRef const *so_buffers,
-				  LLVMValueRef const *so_write_offsets,
-				  struct pipe_stream_output *stream_out,
-				  struct si_shader_output_values *shader_out)
+void si_emit_streamout_output(struct si_shader_context *ctx,
+			      LLVMValueRef const *so_buffers,
+			      LLVMValueRef const *so_write_offsets,
+			      struct pipe_stream_output *stream_out,
+			      struct si_shader_output_values *shader_out)
 {
 	unsigned buf_idx = stream_out->output_buffer;
 	unsigned start = stream_out->start_component;
@@ -2769,13 +2636,16 @@ static void emit_streamout_output(struct si_shader_context *ctx,
 		vdata = out[0];
 		break;
 	case 2: /* as v2i32 */
-	case 3: /* as v4i32 (aligned to 4) */
-	case 4: /* as v4i32 */
-		vdata = LLVMGetUndef(LLVMVectorType(ctx->i32, util_next_power_of_two(num_comps)));
-		for (int j = 0; j < num_comps; j++) {
-			vdata = LLVMBuildInsertElement(ctx->ac.builder, vdata, out[j],
-						       LLVMConstInt(ctx->i32, j, 0), "");
+	case 3: /* as v3i32 */
+		if (ac_has_vec3_support(ctx->screen->info.chip_class, false)) {
+			vdata = ac_build_gather_values(&ctx->ac, out, num_comps);
+			break;
 		}
+		/* as v4i32 (aligned to 4) */
+		out[3] = LLVMGetUndef(ctx->i32);
+		/* fall through */
+	case 4: /* as v4i32 */
+		vdata = ac_build_gather_values(&ctx->ac, out, util_next_power_of_two(num_comps));
 		break;
 	}
 
@@ -2783,7 +2653,7 @@ static void emit_streamout_output(struct si_shader_context *ctx,
 				    vdata, num_comps,
 				    so_write_offsets[buf_idx],
 				    ctx->i32_0,
-				    stream_out->dst_offset * 4, 1, 1, true, false);
+				    stream_out->dst_offset * 4, 1, 1, false);
 }
 
 /**
@@ -2848,9 +2718,9 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
 							      ctx->param_streamout_offset[i]);
 			so_offset = LLVMBuildMul(builder, so_offset, LLVMConstInt(ctx->i32, 4, 0), "");
 
-			so_write_offset[i] = LLVMBuildMul(builder, so_write_index,
-							  LLVMConstInt(ctx->i32, so->stride[i]*4, 0), "");
-			so_write_offset[i] = LLVMBuildAdd(builder, so_write_offset[i], so_offset, "");
+			so_write_offset[i] = ac_build_imad(&ctx->ac, so_write_index,
+							   LLVMConstInt(ctx->i32, so->stride[i]*4, 0),
+							   so_offset);
 		}
 
 		/* Write streamout data. */
@@ -2863,8 +2733,8 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx,
 			if (stream != so->output[i].stream)
 				continue;
 
-			emit_streamout_output(ctx, so_buffers, so_write_offset,
-					      &so->output[i], &outputs[reg]);
+			si_emit_streamout_output(ctx, so_buffers, so_write_offset,
+						 &so->output[i], &outputs[reg]);
 		}
 	}
 	lp_build_endif(&if_ctx);
@@ -2915,7 +2785,8 @@ static void si_build_param_exports(struct si_shader_context *ctx,
 		if ((semantic_name != TGSI_SEMANTIC_GENERIC ||
 		     semantic_index < SI_MAX_IO_GENERIC) &&
 		    shader->key.opt.kill_outputs &
-		    (1ull << si_shader_io_get_unique_index(semantic_name, semantic_index)))
+		    (1ull << si_shader_io_get_unique_index(semantic_name,
+							   semantic_index, true)))
 			continue;
 
 		si_export_param(ctx, param_count, outputs[i].values);
@@ -2927,10 +2798,76 @@ static void si_build_param_exports(struct si_shader_context *ctx,
 	shader->info.nr_param_exports = param_count;
 }
 
-/* Generate export instructions for hardware VS shader stage */
-static void si_llvm_export_vs(struct si_shader_context *ctx,
-			      struct si_shader_output_values *outputs,
-			      unsigned noutput)
+/**
+ * Vertex color clamping.
+ *
+ * This uses a state constant loaded in a user data SGPR and
+ * an IF statement is added that clamps all colors if the constant
+ * is true.
+ */
+static void si_vertex_color_clamping(struct si_shader_context *ctx,
+				     struct si_shader_output_values *outputs,
+				     unsigned noutput)
+{
+	LLVMValueRef addr[SI_MAX_VS_OUTPUTS][4];
+	bool has_colors = false;
+
+	/* Store original colors to alloca variables. */
+	for (unsigned i = 0; i < noutput; i++) {
+		if (outputs[i].semantic_name != TGSI_SEMANTIC_COLOR &&
+		    outputs[i].semantic_name != TGSI_SEMANTIC_BCOLOR)
+			continue;
+
+		for (unsigned j = 0; j < 4; j++) {
+			addr[i][j] = ac_build_alloca_undef(&ctx->ac, ctx->f32, "");
+			LLVMBuildStore(ctx->ac.builder, outputs[i].values[j], addr[i][j]);
+		}
+		has_colors = true;
+	}
+
+	if (!has_colors)
+		return;
+
+	/* The state is in the first bit of the user SGPR. */
+	LLVMValueRef cond = LLVMGetParam(ctx->main_fn, ctx->param_vs_state_bits);
+	cond = LLVMBuildTrunc(ctx->ac.builder, cond, ctx->i1, "");
+
+	struct lp_build_if_state if_ctx;
+	lp_build_if(&if_ctx, &ctx->gallivm, cond);
+
+	/* Store clamped colors to alloca variables within the conditional block. */
+	for (unsigned i = 0; i < noutput; i++) {
+		if (outputs[i].semantic_name != TGSI_SEMANTIC_COLOR &&
+		    outputs[i].semantic_name != TGSI_SEMANTIC_BCOLOR)
+			continue;
+
+		for (unsigned j = 0; j < 4; j++) {
+			LLVMBuildStore(ctx->ac.builder,
+				       ac_build_clamp(&ctx->ac, outputs[i].values[j]),
+				       addr[i][j]);
+		}
+	}
+	lp_build_endif(&if_ctx);
+
+	/* Load clamped colors */
+	for (unsigned i = 0; i < noutput; i++) {
+		if (outputs[i].semantic_name != TGSI_SEMANTIC_COLOR &&
+		    outputs[i].semantic_name != TGSI_SEMANTIC_BCOLOR)
+			continue;
+
+		for (unsigned j = 0; j < 4; j++) {
+			outputs[i].values[j] =
+				LLVMBuildLoad(ctx->ac.builder, addr[i][j], "");
+		}
+	}
+}
+
+/* Generate export instructions for hardware VS shader stage or NGG GS stage
+ * (position and parameter data only).
+ */
+void si_llvm_export_vs(struct si_shader_context *ctx,
+		       struct si_shader_output_values *outputs,
+		       unsigned noutput)
 {
 	struct si_shader *shader = ctx->shader;
 	struct ac_export_args pos_args[4] = {};
@@ -2938,6 +2875,8 @@ static void si_llvm_export_vs(struct si_shader_context *ctx,
 	unsigned pos_idx;
 	int i;
 
+	si_vertex_color_clamping(ctx, outputs, noutput);
+
 	/* Build position exports. */
 	for (i = 0; i < noutput; i++) {
 		switch (outputs[i].semantic_name) {
@@ -3082,7 +3021,7 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
 	LLVMValueRef invocation_id, buffer, buffer_offset;
-	LLVMValueRef lds_vertex_stride, lds_vertex_offset, lds_base;
+	LLVMValueRef lds_vertex_stride, lds_base;
 	uint64_t inputs;
 
 	invocation_id = unpack_llvm_param(ctx, ctx->abi.tcs_rel_ids, 8, 5);
@@ -3090,10 +3029,9 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
 	buffer_offset = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
 
 	lds_vertex_stride = get_tcs_in_vertex_dw_stride(ctx);
-	lds_vertex_offset = LLVMBuildMul(ctx->ac.builder, invocation_id,
-	                                 lds_vertex_stride, "");
 	lds_base = get_tcs_in_current_patch_offset(ctx);
-	lds_base = LLVMBuildAdd(ctx->ac.builder, lds_base, lds_vertex_offset, "");
+	lds_base = ac_build_imad(&ctx->ac, invocation_id, lds_vertex_stride,
+				 lds_base);
 
 	inputs = ctx->shader->key.mono.u.ff_tcs_inputs_to_copy;
 	while (inputs) {
@@ -3108,11 +3046,10 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
 		                              invocation_id,
 		                              LLVMConstInt(ctx->i32, i, 0));
 
-		LLVMValueRef value = lds_load(bld_base, ctx->ac.i32, ~0,
-		                              lds_ptr);
+		LLVMValueRef value = lshs_lds_load(bld_base, ctx->ac.i32, ~0, lds_ptr);
 
 		ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buffer_addr,
-					    buffer_offset, 0, 1, 0, true, false);
+					    buffer_offset, 0, 1, 0, false);
 	}
 }
 
@@ -3195,11 +3132,11 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
 
 		for (i = 0; i < outer_comps; i++) {
 			outer[i] = out[i] =
-				lds_load(bld_base, ctx->ac.i32, i, lds_outer);
+				lshs_lds_load(bld_base, ctx->ac.i32, i, lds_outer);
 		}
 		for (i = 0; i < inner_comps; i++) {
 			inner[i] = out[outer_comps+i] =
-				lds_load(bld_base, ctx->ac.i32, i, lds_inner);
+				lshs_lds_load(bld_base, ctx->ac.i32, i, lds_inner);
 		}
 	}
 
@@ -3213,11 +3150,11 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
 	}
 
 	/* Convert the outputs to vectors for stores. */
-	vec0 = lp_build_gather_values(&ctx->gallivm, out, MIN2(stride, 4));
+	vec0 = ac_build_gather_values(&ctx->ac, out, MIN2(stride, 4));
 	vec1 = NULL;
 
 	if (stride > 4)
-		vec1 = lp_build_gather_values(&ctx->gallivm, out+4, stride - 4);
+		vec1 = ac_build_gather_values(&ctx->ac, out+4, stride - 4);
 
 	/* Get the buffer. */
 	buffer = get_tess_ring_descriptor(ctx, TCS_FACTOR_RING);
@@ -3234,11 +3171,11 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
 
 	/* Store the dynamic HS control word. */
 	offset = 0;
-	if (ctx->screen->info.chip_class <= VI) {
+	if (ctx->screen->info.chip_class <= GFX8) {
 		ac_build_buffer_store_dword(&ctx->ac, buffer,
 					    LLVMConstInt(ctx->i32, 0x80000000, 0),
 					    1, ctx->i32_0, tf_base,
-					    offset, 1, 0, true, false);
+					    offset, 1, 0, false);
 		offset += 4;
 	}
 
@@ -3247,12 +3184,12 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
 	/* Store the tessellation factors. */
 	ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
 				    MIN2(stride, 4), byteoffset, tf_base,
-				    offset, 1, 0, true, false);
+				    offset, 1, 0, false);
 	offset += 16;
 	if (vec1)
 		ac_build_buffer_store_dword(&ctx->ac, buffer, vec1,
 					    stride - 4, byteoffset, tf_base,
-					    offset, 1, 0, true, false);
+					    offset, 1, 0, false);
 
 	/* Store the tess factors into the offchip buffer if TES reads them. */
 	if (shader->key.part.tcs.epilog.tes_reads_tess_factors) {
@@ -3268,12 +3205,14 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
 		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(&ctx->gallivm, outer,
-						   util_next_power_of_two(outer_comps));
+		unsigned outer_vec_size =
+			ac_has_vec3_support(ctx->screen->info.chip_class, false) ?
+				outer_comps : util_next_power_of_two(outer_comps);
+		outer_vec = ac_build_gather_values(&ctx->ac, outer, outer_vec_size);
 
 		ac_build_buffer_store_dword(&ctx->ac, buf, outer_vec,
 					    outer_comps, tf_outer_offset,
-					    base, 0, 1, 0, true, false);
+					    base, 0, 1, 0, false);
 		if (inner_comps) {
 			param_inner = si_shader_io_get_unique_index_patch(
 					      TGSI_SEMANTIC_TESSINNER, 0);
@@ -3281,10 +3220,10 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
 					LLVMConstInt(ctx->i32, param_inner, 0));
 
 			inner_vec = inner_comps == 1 ? inner[0] :
-				    lp_build_gather_values(&ctx->gallivm, inner, inner_comps);
+				    ac_build_gather_values(&ctx->ac, inner, inner_comps);
 			ac_build_buffer_store_dword(&ctx->ac, buf, inner_vec,
 						    inner_comps, tf_inner_offset,
-						    base, 0, 1, 0, true, false);
+						    base, 0, 1, 0, false);
 		}
 	}
 
@@ -3317,21 +3256,9 @@ si_insert_input_ptr(struct si_shader_context *ctx, LLVMValueRef ret,
 		    unsigned param, unsigned return_index)
 {
 	LLVMBuilderRef builder = ctx->ac.builder;
-	LLVMValueRef ptr, lo, hi;
-
-	if (HAVE_32BIT_POINTERS) {
-		ptr = LLVMGetParam(ctx->main_fn, param);
-		ptr = LLVMBuildPtrToInt(builder, ptr, ctx->i32, "");
-		return LLVMBuildInsertValue(builder, ret, ptr, return_index, "");
-	}
-
-	ptr = LLVMGetParam(ctx->main_fn, param);
-	ptr = LLVMBuildPtrToInt(builder, ptr, ctx->i64, "");
-	ptr = LLVMBuildBitCast(builder, ptr, ctx->v2i32, "");
-	lo = LLVMBuildExtractElement(builder, ptr, ctx->i32_0, "");
-	hi = LLVMBuildExtractElement(builder, ptr, ctx->i32_1, "");
-	ret = LLVMBuildInsertValue(builder, ret, lo, return_index, "");
-	return LLVMBuildInsertValue(builder, ret, hi, return_index + 1, "");
+	LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, param);
+	ptr = LLVMBuildPtrToInt(builder, ptr, ctx->i32, "");
+	return LLVMBuildInsertValue(builder, ret, ptr, return_index, "");
 }
 
 /* This only writes the tessellation factor levels. */
@@ -3432,8 +3359,7 @@ static void si_set_ls_return_value_for_tcs(struct si_shader_context *ctx)
 	LLVMValueRef ret = ctx->return_value;
 
 	ret = si_insert_input_ptr(ctx, ret, 0, 0);
-	if (HAVE_32BIT_POINTERS)
-		ret = si_insert_input_ptr(ctx, ret, 1, 1);
+	ret = si_insert_input_ptr(ctx, ret, 1, 1);
 	ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_offchip_offset, 2);
 	ret = si_insert_input_ret(ctx, ret, ctx->param_merged_wave_info, 3);
 	ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_factor_offset, 4);
@@ -3448,11 +3374,6 @@ static void si_set_ls_return_value_for_tcs(struct si_shader_context *ctx)
 	ret = si_insert_input_ret(ctx, ret, ctx->param_vs_state_bits,
 				  8 + SI_SGPR_VS_STATE_BITS);
 
-#if !HAVE_32BIT_POINTERS
-	ret = si_insert_input_ptr(ctx, ret, ctx->param_vs_state_bits + 1,
-				  8 + GFX9_SGPR_2ND_SAMPLERS_AND_IMAGES);
-#endif
-
 	ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_offchip_layout,
 				  8 + GFX9_SGPR_TCS_OFFCHIP_LAYOUT);
 	ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_out_lds_offsets,
@@ -3476,8 +3397,7 @@ static void si_set_es_return_value_for_gs(struct si_shader_context *ctx)
 	LLVMValueRef ret = ctx->return_value;
 
 	ret = si_insert_input_ptr(ctx, ret, 0, 0);
-	if (HAVE_32BIT_POINTERS)
-		ret = si_insert_input_ptr(ctx, ret, 1, 1);
+	ret = si_insert_input_ptr(ctx, ret, 1, 1);
 	ret = si_insert_input_ret(ctx, ret, ctx->param_gs2vs_offset, 2);
 	ret = si_insert_input_ret(ctx, ret, ctx->param_merged_wave_info, 3);
 	ret = si_insert_input_ret(ctx, ret, ctx->param_merged_scratch_offset, 5);
@@ -3488,11 +3408,6 @@ static void si_set_es_return_value_for_gs(struct si_shader_context *ctx)
 				  ctx->param_bindless_samplers_and_images,
 				  8 + SI_SGPR_BINDLESS_SAMPLERS_AND_IMAGES);
 
-#if !HAVE_32BIT_POINTERS
-	ret = si_insert_input_ptr(ctx, ret, ctx->param_vs_state_bits + 1,
-				  8 + GFX9_SGPR_2ND_SAMPLERS_AND_IMAGES);
-#endif
-
 	unsigned vgpr;
 	if (ctx->type == PIPE_SHADER_VERTEX)
 		vgpr = 8 + GFX9_VSGS_NUM_USER_SGPR;
@@ -3545,7 +3460,7 @@ static void si_llvm_emit_ls_epilogue(struct ac_shader_abi *abi,
 		    name == TGSI_SEMANTIC_VIEWPORT_INDEX)
 			continue;
 
-		int param = si_shader_io_get_unique_index(name, index);
+		int param = si_shader_io_get_unique_index(name, index, false);
 		LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr,
 					LLVMConstInt(ctx->i32, param * 4, 0), "");
 
@@ -3553,7 +3468,7 @@ static void si_llvm_emit_ls_epilogue(struct ac_shader_abi *abi,
 			if (!(info->output_usagemask[i] & (1 << chan)))
 				continue;
 
-			lds_store(ctx, chan, dw_addr,
+			lshs_lds_store(ctx, chan, dw_addr,
 				  LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], ""));
 		}
 	}
@@ -3594,15 +3509,20 @@ static void si_llvm_emit_es_epilogue(struct ac_shader_abi *abi,
 			continue;
 
 		param = si_shader_io_get_unique_index(info->output_semantic_name[i],
-						      info->output_semantic_index[i]);
+						      info->output_semantic_index[i], false);
 
 		for (chan = 0; chan < 4; chan++) {
+			if (!(info->output_usagemask[i] & (1 << chan)))
+				continue;
+
 			LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], "");
 			out_val = ac_to_integer(&ctx->ac, out_val);
 
 			/* GFX9 has the ESGS ring in LDS. */
 			if (ctx->screen->info.chip_class >= GFX9) {
-				lds_store(ctx, param * 4 + chan, lds_base, out_val);
+				LLVMValueRef idx = LLVMConstInt(ctx->i32, param * 4 + chan, false);
+				idx = LLVMBuildAdd(ctx->ac.builder, lds_base, idx, "");
+				ac_build_indexed_store(&ctx->ac, ctx->esgs_ring, idx, out_val);
 				continue;
 			}
 
@@ -3610,7 +3530,7 @@ static void si_llvm_emit_es_epilogue(struct ac_shader_abi *abi,
 						    ctx->esgs_ring,
 						    out_val, 1, NULL, soffset,
 						    (4 * param + chan) * 4,
-						    1, 1, true, true);
+						    1, 1, true);
 		}
 	}
 
@@ -3667,44 +3587,6 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi,
 
 	outputs = MALLOC((info->num_outputs + 1) * sizeof(outputs[0]));
 
-	/* Vertex color clamping.
-	 *
-	 * This uses a state constant loaded in a user data SGPR and
-	 * an IF statement is added that clamps all colors if the constant
-	 * is true.
-	 */
-	if (ctx->type == PIPE_SHADER_VERTEX) {
-		struct lp_build_if_state if_ctx;
-		LLVMValueRef cond = NULL;
-		LLVMValueRef addr, val;
-
-		for (i = 0; i < info->num_outputs; i++) {
-			if (info->output_semantic_name[i] != TGSI_SEMANTIC_COLOR &&
-			    info->output_semantic_name[i] != TGSI_SEMANTIC_BCOLOR)
-				continue;
-
-			/* We've found a color. */
-			if (!cond) {
-				/* The state is in the first bit of the user SGPR. */
-				cond = LLVMGetParam(ctx->main_fn,
-						    ctx->param_vs_state_bits);
-				cond = LLVMBuildTrunc(ctx->ac.builder, cond,
-						      ctx->i1, "");
-				lp_build_if(&if_ctx, &ctx->gallivm, cond);
-			}
-
-			for (j = 0; j < 4; j++) {
-				addr = addrs[4 * i + j];
-				val = LLVMBuildLoad(ctx->ac.builder, addr, "");
-				val = ac_build_clamp(&ctx->ac, val);
-				LLVMBuildStore(ctx->ac.builder, val, addr);
-			}
-		}
-
-		if (cond)
-			lp_build_endif(&if_ctx);
-	}
-
 	for (i = 0; i < info->num_outputs; i++) {
 		outputs[i].semantic_name = info->output_semantic_name[i];
 		outputs[i].semantic_index = info->output_semantic_index[i];
@@ -3726,7 +3608,7 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi,
 	if (ctx->shader->key.mono.u.vs_export_prim_id) {
 		outputs[i].semantic_name = TGSI_SEMANTIC_PRIMID;
 		outputs[i].semantic_index = 0;
-		outputs[i].values[0] = ac_to_float(&ctx->ac, get_primitive_id(ctx, 0));
+		outputs[i].values[0] = ac_to_float(&ctx->ac, si_get_primitive_id(ctx, 0));
 		for (j = 1; j < 4; j++)
 			outputs[i].values[j] = LLVMConstReal(ctx->f32, 0);
 
@@ -3739,6 +3621,33 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi,
 	FREE(outputs);
 }
 
+static void si_llvm_emit_prim_discard_cs_epilogue(struct ac_shader_abi *abi,
+						  unsigned max_outputs,
+						  LLVMValueRef *addrs)
+{
+	struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+	struct tgsi_shader_info *info = &ctx->shader->selector->info;
+	LLVMValueRef pos[4] = {};
+
+	assert(info->num_outputs <= max_outputs);
+
+	for (unsigned i = 0; i < info->num_outputs; i++) {
+		if (info->output_semantic_name[i] != TGSI_SEMANTIC_POSITION)
+			continue;
+
+		for (unsigned chan = 0; chan < 4; chan++)
+			pos[chan] = LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], "");
+		break;
+	}
+	assert(pos[0] != NULL);
+
+	/* Return the position output. */
+	LLVMValueRef ret = ctx->return_value;
+	for (unsigned chan = 0; chan < 4; chan++)
+		ret = LLVMBuildInsertValue(ctx->ac.builder, ret, pos[chan], chan, "");
+	ctx->return_value = ret;
+}
+
 static void si_tgsi_emit_epilogue(struct lp_build_tgsi_context *bld_base)
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
@@ -3894,7 +3803,7 @@ static void si_llvm_return_fs_outputs(struct ac_shader_abi *abi,
 						   addrs[4 * i + 0], "");
 			break;
 		default:
-			fprintf(stderr, "Warning: SI unhandled fs output type:%d\n",
+			fprintf(stderr, "Warning: GFX6 unhandled fs output type:%d\n",
 				semantic_name);
 		}
 	}
@@ -4000,45 +3909,31 @@ static void si_llvm_emit_ddxy(
 	emit_data->output[emit_data->chan] = val;
 }
 
-/*
- * this takes an I,J coordinate pair,
- * and works out the X and Y derivatives.
- * it returns DDX(I), DDX(J), DDY(I), DDY(J).
- */
-static LLVMValueRef si_llvm_emit_ddxy_interp(
-	struct lp_build_tgsi_context *bld_base,
-	LLVMValueRef interp_ij)
-{
-	struct si_shader_context *ctx = si_shader_context(bld_base);
-	LLVMValueRef result[4], a;
-	unsigned i;
-
-	for (i = 0; i < 2; i++) {
-		a = LLVMBuildExtractElement(ctx->ac.builder, interp_ij,
-					    LLVMConstInt(ctx->i32, i, 0), "");
-		result[i] = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_DDX, a);
-		result[2+i] = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_DDY, a);
-	}
-
-	return lp_build_gather_values(&ctx->gallivm, result, 4);
-}
-
-static void interp_fetch_args(
-	struct lp_build_tgsi_context *bld_base,
-	struct lp_build_emit_data *emit_data)
+static void build_interp_intrinsic(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 si_shader *shader = ctx->shader;
+	const struct tgsi_shader_info *info = &shader->selector->info;
+	LLVMValueRef interp_param;
 	const struct tgsi_full_instruction *inst = emit_data->inst;
+	const struct tgsi_full_src_register *input = &inst->Src[0];
+	int input_base, input_array_size;
+	int chan;
+	int i;
+	LLVMValueRef prim_mask = ctx->abi.prim_mask;
+	LLVMValueRef array_idx, offset_x = NULL, offset_y = NULL;
+	int interp_param_idx;
+	unsigned interp;
+	unsigned location;
 
 	if (inst->Instruction.Opcode == TGSI_OPCODE_INTERP_OFFSET) {
 		/* offset is in second src, first two channels */
-		emit_data->args[0] = lp_build_emit_fetch(bld_base,
-							 emit_data->inst, 1,
-							 TGSI_CHAN_X);
-		emit_data->args[1] = lp_build_emit_fetch(bld_base,
-							 emit_data->inst, 1,
-							 TGSI_CHAN_Y);
-		emit_data->arg_count = 2;
+		offset_x = lp_build_emit_fetch(bld_base, emit_data->inst, 1,
+					       TGSI_CHAN_X);
+		offset_y = lp_build_emit_fetch(bld_base, emit_data->inst, 1,
+					       TGSI_CHAN_Y);
 	} else if (inst->Instruction.Opcode == TGSI_OPCODE_INTERP_SAMPLE) {
 		LLVMValueRef sample_position;
 		LLVMValueRef sample_id;
@@ -4073,42 +3968,19 @@ static void interp_fetch_args(
 				ctx->ac.f32_0,
 			};
 
-			sample_position = lp_build_gather_values(&ctx->gallivm, center, 4);
+			sample_position = ac_build_gather_values(&ctx->ac, center, 4);
 		} else {
 			sample_position = load_sample_position(&ctx->abi, sample_id);
 		}
 
-		emit_data->args[0] = LLVMBuildExtractElement(ctx->ac.builder,
-							     sample_position,
-							     ctx->i32_0, "");
+		offset_x = LLVMBuildExtractElement(ctx->ac.builder, sample_position,
+						   ctx->i32_0, "");
 
-		emit_data->args[0] = LLVMBuildFSub(ctx->ac.builder, emit_data->args[0], halfval, "");
-		emit_data->args[1] = LLVMBuildExtractElement(ctx->ac.builder,
-							     sample_position,
-							     ctx->i32_1, "");
-		emit_data->args[1] = LLVMBuildFSub(ctx->ac.builder, emit_data->args[1], halfval, "");
-		emit_data->arg_count = 2;
+		offset_x = LLVMBuildFSub(ctx->ac.builder, offset_x, halfval, "");
+		offset_y = LLVMBuildExtractElement(ctx->ac.builder, sample_position,
+						   ctx->i32_1, "");
+		offset_y = LLVMBuildFSub(ctx->ac.builder, offset_y, halfval, "");
 	}
-}
-
-static void build_interp_intrinsic(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 si_shader *shader = ctx->shader;
-	const struct tgsi_shader_info *info = &shader->selector->info;
-	LLVMValueRef interp_param;
-	const struct tgsi_full_instruction *inst = emit_data->inst;
-	const struct tgsi_full_src_register *input = &inst->Src[0];
-	int input_base, input_array_size;
-	int chan;
-	int i;
-	LLVMValueRef prim_mask = ctx->abi.prim_mask;
-	LLVMValueRef array_idx;
-	int interp_param_idx;
-	unsigned interp;
-	unsigned location;
 
 	assert(input->Register.File == TGSI_FILE_INPUT);
 
@@ -4150,7 +4022,7 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
 	if (inst->Instruction.Opcode == TGSI_OPCODE_INTERP_OFFSET ||
 	    inst->Instruction.Opcode == TGSI_OPCODE_INTERP_SAMPLE) {
 		LLVMValueRef ij_out[2];
-		LLVMValueRef ddxy_out = si_llvm_emit_ddxy_interp(bld_base, interp_param);
+		LLVMValueRef ddxy_out = ac_build_ddxy_interp(&ctx->ac, interp_param);
 
 		/*
 		 * take the I then J parameters, and the DDX/Y for it, and
@@ -4169,19 +4041,14 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
 								      ddxy_out, iy_ll, "");
 			LLVMValueRef interp_el = LLVMBuildExtractElement(ctx->ac.builder,
 									 interp_param, ix_ll, "");
-			LLVMValueRef temp1, temp2;
+			LLVMValueRef temp;
 
 			interp_el = ac_to_float(&ctx->ac, interp_el);
 
-			temp1 = LLVMBuildFMul(ctx->ac.builder, ddx_el, emit_data->args[0], "");
-
-			temp1 = LLVMBuildFAdd(ctx->ac.builder, temp1, interp_el, "");
-
-			temp2 = LLVMBuildFMul(ctx->ac.builder, ddy_el, emit_data->args[1], "");
-
-			ij_out[i] = LLVMBuildFAdd(ctx->ac.builder, temp2, temp1, "");
+			temp = ac_build_fmad(&ctx->ac, ddx_el, offset_x, interp_el);
+			ij_out[i] = ac_build_fmad(&ctx->ac, ddy_el, offset_y, temp);
 		}
-		interp_param = lp_build_gather_values(&ctx->gallivm, ij_out, 2);
+		interp_param = ac_build_gather_values(&ctx->ac, ij_out, 2);
 	}
 
 	if (interp_param)
@@ -4265,19 +4132,6 @@ static void ballot_emit(
 	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,
@@ -4285,6 +4139,16 @@ static void read_lane_emit(
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
 
+	if (emit_data->inst->Instruction.Opcode == TGSI_OPCODE_READ_INVOC) {
+		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;
+	}
+
 	/* We currently have no other way to prevent LLVM from lifting the icmp
 	 * calls to a dominating basic block.
 	 */
@@ -4322,7 +4186,6 @@ static void si_llvm_emit_vertex(struct ac_shader_abi *abi,
 {
 	struct si_shader_context *ctx = si_shader_context_from_abi(abi);
 	struct tgsi_shader_info *info = &ctx->shader->selector->info;
-	struct lp_build_context *uint = &ctx->bld_base.uint_bld;
 	struct si_shader *shader = ctx->shader;
 	struct lp_build_if_state if_state;
 	LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
@@ -4369,8 +4232,9 @@ static void si_llvm_emit_vertex(struct ac_shader_abi *abi,
 					     shader->selector->gs_max_out_vertices, 0);
 			offset++;
 
-			voffset = lp_build_add(uint, voffset, gs_next_vertex);
-			voffset = lp_build_mul_imm(uint, voffset, 4);
+			voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
+			voffset = LLVMBuildMul(ctx->ac.builder, voffset,
+					       LLVMConstInt(ctx->i32, 4, 0), "");
 
 			out_val = ac_to_integer(&ctx->ac, out_val);
 
@@ -4378,18 +4242,19 @@ static void si_llvm_emit_vertex(struct ac_shader_abi *abi,
 						    ctx->gsvs_ring[stream],
 						    out_val, 1,
 						    voffset, soffset, 0,
-						    1, 1, true, true);
+						    1, 1, true);
 		}
 	}
 
-	gs_next_vertex = lp_build_add(uint, gs_next_vertex,
-				      ctx->i32_1);
-
+	gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex, ctx->i32_1, "");
 	LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
 
-	/* Signal vertex emission */
-	ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
-			 si_get_gs_wave_id(ctx));
+	/* Signal vertex emission if vertex data was written. */
+	if (offset) {
+		ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
+				 si_get_gs_wave_id(ctx));
+	}
+
 	if (!use_kill)
 		lp_build_endif(&if_state);
 }
@@ -4434,31 +4299,24 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
 
-	/* SI only (thanks to a hw bug workaround):
+	/* GFX6 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->screen->info.chip_class == SI &&
+	if (ctx->screen->info.chip_class == GFX6 &&
 	    ctx->type == PIPE_SHADER_TESS_CTRL) {
 		ac_build_waitcnt(&ctx->ac, LGKM_CNT & VM_CNT);
 		return;
 	}
 
-	lp_build_intrinsic(ctx->ac.builder,
-			   "llvm.amdgcn.s.barrier",
-			   ctx->voidt, NULL, 0, LP_FUNC_ATTR_CONVERGENT);
+	ac_build_s_barrier(&ctx->ac);
 }
 
-static const struct lp_build_tgsi_action interp_action = {
-	.fetch_args = interp_fetch_args,
-	.emit = build_interp_intrinsic,
-};
-
-static void si_create_function(struct si_shader_context *ctx,
-			       const char *name,
-			       LLVMTypeRef *returns, unsigned num_returns,
-			       struct si_function_info *fninfo,
-			       unsigned max_workgroup_size)
+void si_create_function(struct si_shader_context *ctx,
+			const char *name,
+			LLVMTypeRef *returns, unsigned num_returns,
+			struct si_function_info *fninfo,
+			unsigned max_workgroup_size)
 {
 	int i;
 
@@ -4476,10 +4334,12 @@ static void si_create_function(struct si_shader_context *ctx,
 		 * allows the optimization passes to move loads and reduces
 		 * SGPR spilling significantly.
 		 */
-		lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_INREG);
+		ac_add_function_attr(ctx->ac.context, ctx->main_fn, i + 1,
+				     AC_FUNC_ATTR_INREG);
 
 		if (LLVMGetTypeKind(LLVMTypeOf(P)) == LLVMPointerTypeKind) {
-			lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_NOALIAS);
+			ac_add_function_attr(ctx->ac.context, ctx->main_fn, i + 1,
+					     AC_FUNC_ATTR_NOALIAS);
 			ac_add_attr_dereferenceable(P, UINT64_MAX);
 		}
 	}
@@ -4495,11 +4355,8 @@ static void si_create_function(struct si_shader_context *ctx,
 						     ctx->screen->info.address32_hi);
 	}
 
-	if (max_workgroup_size) {
-		ac_llvm_add_target_dep_function_attr(ctx->main_fn,
-						     "amdgpu-max-work-group-size",
-						     max_workgroup_size);
-	}
+	ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
+
 	LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
 					   "no-signed-zeros-fp-math",
 					   "true");
@@ -4551,7 +4408,7 @@ static unsigned si_get_max_workgroup_size(const struct si_shader *shader)
 	case PIPE_SHADER_TESS_CTRL:
 		/* Return this so that LLVM doesn't remove s_barrier
 		 * instructions on chips where we use s_barrier. */
-		return shader->selector->screen->info.chip_class >= CIK ? 128 : 64;
+		return shader->selector->screen->info.chip_class >= GFX7 ? 128 : 64;
 
 	case PIPE_SHADER_GEOMETRY:
 		return shader->selector->screen->info.chip_class >= GFX9 ? 128 : 64;
@@ -4630,10 +4487,10 @@ static void declare_global_desc_pointers(struct si_shader_context *ctx,
 static void declare_vs_specific_input_sgprs(struct si_shader_context *ctx,
 					    struct si_function_info *fninfo)
 {
+	ctx->param_vs_state_bits = add_arg(fninfo, ARG_SGPR, ctx->i32);
 	add_arg_assign(fninfo, ARG_SGPR, ctx->i32, &ctx->abi.base_vertex);
 	add_arg_assign(fninfo, ARG_SGPR, ctx->i32, &ctx->abi.start_instance);
 	add_arg_assign(fninfo, ARG_SGPR, ctx->i32, &ctx->abi.draw_id);
-	ctx->param_vs_state_bits = add_arg(fninfo, ARG_SGPR, ctx->i32);
 }
 
 static void declare_vs_input_vgprs(struct si_shader_context *ctx,
@@ -4661,6 +4518,30 @@ static void declare_vs_input_vgprs(struct si_shader_context *ctx,
 	}
 }
 
+static void declare_vs_blit_inputs(struct si_shader_context *ctx,
+				   struct si_function_info *fninfo,
+				   unsigned vs_blit_property)
+{
+	ctx->param_vs_blit_inputs = fninfo->num_params;
+	add_arg(fninfo, ARG_SGPR, ctx->i32); /* i16 x1, y1 */
+	add_arg(fninfo, ARG_SGPR, ctx->i32); /* i16 x2, y2 */
+	add_arg(fninfo, ARG_SGPR, ctx->f32); /* depth */
+
+	if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_COLOR) {
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* color0 */
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* color1 */
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* color2 */
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* color3 */
+	} else if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_TEXCOORD) {
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* texcoord.x1 */
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* texcoord.y1 */
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* texcoord.x2 */
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* texcoord.y2 */
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* texcoord.z */
+		add_arg(fninfo, ARG_SGPR, ctx->f32); /* texcoord.w */
+	}
+}
+
 static void declare_tes_input_vgprs(struct si_shader_context *ctx,
 				    struct si_function_info *fninfo)
 {
@@ -4705,24 +4586,7 @@ static void create_function(struct si_shader_context *ctx)
 		declare_global_desc_pointers(ctx, &fninfo);
 
 		if (vs_blit_property) {
-			ctx->param_vs_blit_inputs = fninfo.num_params;
-			add_arg(&fninfo, ARG_SGPR, ctx->i32); /* i16 x1, y1 */
-			add_arg(&fninfo, ARG_SGPR, ctx->i32); /* i16 x2, y2 */
-			add_arg(&fninfo, ARG_SGPR, ctx->f32); /* depth */
-
-			if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_COLOR) {
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* color0 */
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* color1 */
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* color2 */
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* color3 */
-			} else if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_TEXCOORD) {
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.x1 */
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.y1 */
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.x2 */
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.y2 */
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.z */
-				add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.w */
-			}
+			declare_vs_blit_inputs(ctx, &fninfo, vs_blit_property);
 
 			/* VGPRs */
 			declare_vs_input_vgprs(ctx, &fninfo, &num_prolog_vgprs);
@@ -4740,7 +4604,7 @@ static void create_function(struct si_shader_context *ctx)
 			/* no extra parameters */
 		} else {
 			if (shader->is_gs_copy_shader) {
-				fninfo.num_params = ctx->param_rw_buffers + 1;
+				fninfo.num_params = ctx->param_vs_state_bits + 1;
 				fninfo.num_sgpr_params = fninfo.num_params;
 			}
 
@@ -4751,9 +4615,15 @@ static void create_function(struct si_shader_context *ctx)
 
 		/* VGPRs */
 		declare_vs_input_vgprs(ctx, &fninfo, &num_prolog_vgprs);
+
+		/* Return values */
+		if (shader->key.opt.vs_as_prim_discard_cs) {
+			for (i = 0; i < 4; i++)
+				returns[num_returns++] = ctx->f32; /* VGPRs */
+		}
 		break;
 
-	case PIPE_SHADER_TESS_CTRL: /* SI-CI-VI */
+	case PIPE_SHADER_TESS_CTRL: /* GFX6-GFX8 */
 		declare_global_desc_pointers(ctx, &fninfo);
 		declare_per_stage_desc_pointers(ctx, &fninfo, true);
 		ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
@@ -4779,13 +4649,8 @@ static void create_function(struct si_shader_context *ctx)
 	case SI_SHADER_MERGED_VERTEX_TESSCTRL:
 		/* Merged stages have 8 system SGPRs at the beginning. */
 		/* SPI_SHADER_USER_DATA_ADDR_LO/HI_HS */
-		if (HAVE_32BIT_POINTERS) {
-			declare_per_stage_desc_pointers(ctx, &fninfo,
-							ctx->type == PIPE_SHADER_TESS_CTRL);
-		} else {
-			declare_const_and_shader_buffers(ctx, &fninfo,
-							 ctx->type == PIPE_SHADER_TESS_CTRL);
-		}
+		declare_per_stage_desc_pointers(ctx, &fninfo,
+						ctx->type == PIPE_SHADER_TESS_CTRL);
 		ctx->param_tcs_offchip_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_merged_wave_info = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_tcs_factor_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
@@ -4798,15 +4663,9 @@ static void create_function(struct si_shader_context *ctx)
 						ctx->type == PIPE_SHADER_VERTEX);
 		declare_vs_specific_input_sgprs(ctx, &fninfo);
 
-		if (!HAVE_32BIT_POINTERS) {
-			declare_samplers_and_images(ctx, &fninfo,
-						    ctx->type == PIPE_SHADER_TESS_CTRL);
-		}
 		ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_tcs_out_lds_offsets = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_tcs_out_lds_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
-		if (!HAVE_32BIT_POINTERS) /* Align to 2 dwords. */
-			add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
 		ctx->param_vertex_buffers = add_arg(&fninfo, ARG_SGPR,
 			ac_array_in_const32_addr_space(ctx->v4i32));
 
@@ -4840,13 +4699,8 @@ static void create_function(struct si_shader_context *ctx)
 	case SI_SHADER_MERGED_VERTEX_OR_TESSEVAL_GEOMETRY:
 		/* Merged stages have 8 system SGPRs at the beginning. */
 		/* SPI_SHADER_USER_DATA_ADDR_LO/HI_GS */
-		if (HAVE_32BIT_POINTERS) {
-			declare_per_stage_desc_pointers(ctx, &fninfo,
-							ctx->type == PIPE_SHADER_GEOMETRY);
-		} else {
-			declare_const_and_shader_buffers(ctx, &fninfo,
-							 ctx->type == PIPE_SHADER_GEOMETRY);
-		}
+		declare_per_stage_desc_pointers(ctx, &fninfo,
+						ctx->type == PIPE_SHADER_GEOMETRY);
 		ctx->param_gs2vs_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_merged_wave_info = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_tcs_offchip_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
@@ -4861,19 +4715,12 @@ static void create_function(struct si_shader_context *ctx)
 		if (ctx->type == PIPE_SHADER_VERTEX) {
 			declare_vs_specific_input_sgprs(ctx, &fninfo);
 		} else {
+			ctx->param_vs_state_bits = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 			ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 			ctx->param_tes_offchip_addr = add_arg(&fninfo, ARG_SGPR, ctx->i32);
-			if (!HAVE_32BIT_POINTERS) {
-				/* Declare as many input SGPRs as the VS has. */
-				add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
-				ctx->param_vs_state_bits = add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
-			}
+			/* Declare as many input SGPRs as the VS has. */
 		}
 
-		if (!HAVE_32BIT_POINTERS) {
-			declare_samplers_and_images(ctx, &fninfo,
-						    ctx->type == PIPE_SHADER_GEOMETRY);
-		}
 		if (ctx->type == PIPE_SHADER_VERTEX) {
 			ctx->param_vertex_buffers = add_arg(&fninfo, ARG_SGPR,
 				ac_array_in_const32_addr_space(ctx->v4i32));
@@ -4913,6 +4760,7 @@ static void create_function(struct si_shader_context *ctx)
 	case PIPE_SHADER_TESS_EVAL:
 		declare_global_desc_pointers(ctx, &fninfo);
 		declare_per_stage_desc_pointers(ctx, &fninfo, true);
+		ctx->param_vs_state_bits = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_tes_offchip_addr = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 
@@ -5018,9 +4866,17 @@ static void create_function(struct si_shader_context *ctx)
 		declare_per_stage_desc_pointers(ctx, &fninfo, true);
 		if (shader->selector->info.uses_grid_size)
 			add_arg_assign(&fninfo, ARG_SGPR, v3i32, &ctx->abi.num_work_groups);
-		if (shader->selector->info.uses_block_size)
+		if (shader->selector->info.uses_block_size &&
+		    shader->selector->info.properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0)
 			ctx->param_block_size = add_arg(&fninfo, ARG_SGPR, v3i32);
 
+		unsigned cs_user_data_dwords =
+			shader->selector->info.properties[TGSI_PROPERTY_CS_USER_DATA_DWORDS];
+		if (cs_user_data_dwords) {
+			ctx->param_cs_user_data = add_arg(&fninfo, ARG_SGPR,
+							  LLVMVectorType(ctx->i32, cs_user_data_dwords));
+		}
+
 		for (i = 0; i < 3; i++) {
 			ctx->abi.workgroup_ids[i] = NULL;
 			if (shader->selector->info.uses_block_id[i])
@@ -5038,8 +4894,7 @@ static void create_function(struct si_shader_context *ctx)
 			   si_get_max_workgroup_size(shader));
 
 	/* Reserve register locations for VGPR inputs the PS prolog may need. */
-	if (ctx->type == PIPE_SHADER_FRAGMENT &&
-	    ctx->separate_prolog) {
+	if (ctx->type == PIPE_SHADER_FRAGMENT && !ctx->shader->is_monolithic) {
 		ac_llvm_add_target_dep_function_attr(ctx->main_fn,
 						     "InitialPSInputAddr",
 						     S_0286D0_PERSP_SAMPLE_ENA(1) |
@@ -5065,11 +4920,21 @@ static void create_function(struct si_shader_context *ctx)
 	assert(shader->info.num_input_vgprs >= num_prolog_vgprs);
 	shader->info.num_input_vgprs -= num_prolog_vgprs;
 
-	if (shader->key.as_ls ||
-	    ctx->type == PIPE_SHADER_TESS_CTRL ||
-	    /* GFX9 has the ESGS ring buffer in LDS. */
-	    type == SI_SHADER_MERGED_VERTEX_OR_TESSEVAL_GEOMETRY)
-		ac_declare_lds_as_pointer(&ctx->ac);
+	if (shader->key.as_ls || ctx->type == PIPE_SHADER_TESS_CTRL) {
+		if (USE_LDS_SYMBOLS && HAVE_LLVM >= 0x0900) {
+			/* The LSHS size is not known until draw time, so we append it
+			 * at the end of whatever LDS use there may be in the rest of
+			 * the shader (currently none, unless LLVM decides to do its
+			 * own LDS-based lowering).
+			 */
+			ctx->ac.lds = LLVMAddGlobalInAddressSpace(
+				ctx->ac.module, LLVMArrayType(ctx->i32, 0),
+				"__lds_end", AC_ADDR_SPACE_LDS);
+			LLVMSetAlignment(ctx->ac.lds, 256);
+		} else {
+			ac_declare_lds_as_pointer(&ctx->ac);
+		}
+	}
 }
 
 /**
@@ -5083,15 +4948,33 @@ static void preload_ring_buffers(struct si_shader_context *ctx)
 	LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
 					    ctx->param_rw_buffers);
 
-	if (ctx->screen->info.chip_class <= VI &&
-	    (ctx->shader->key.as_es || ctx->type == PIPE_SHADER_GEOMETRY)) {
-		unsigned ring =
-			ctx->type == PIPE_SHADER_GEOMETRY ? SI_GS_RING_ESGS
-							     : SI_ES_RING_ESGS;
-		LLVMValueRef offset = LLVMConstInt(ctx->i32, ring, 0);
+	if (ctx->shader->key.as_es || ctx->type == PIPE_SHADER_GEOMETRY) {
+		if (ctx->screen->info.chip_class <= GFX8) {
+			unsigned ring =
+				ctx->type == PIPE_SHADER_GEOMETRY ? SI_GS_RING_ESGS
+								  : SI_ES_RING_ESGS;
+			LLVMValueRef offset = LLVMConstInt(ctx->i32, ring, 0);
 
-		ctx->esgs_ring =
-			ac_build_load_to_sgpr(&ctx->ac, buf_ptr, offset);
+			ctx->esgs_ring =
+				ac_build_load_to_sgpr(&ctx->ac, buf_ptr, offset);
+		} else {
+			if (USE_LDS_SYMBOLS && HAVE_LLVM >= 0x0900) {
+				/* Declare the ESGS ring as an explicit LDS symbol.
+				 * For monolithic shaders, we declare the ring only once.
+				 *
+				 * We declare it with 64KB alignment as a hint that the
+				 * pointer value will always be 0.
+				 */
+				ctx->esgs_ring = LLVMAddGlobalInAddressSpace(
+					ctx->ac.module, LLVMArrayType(ctx->i32, 0),
+					"esgs_ring",
+					AC_ADDR_SPACE_LDS);
+				LLVMSetAlignment(ctx->esgs_ring, 64 * 1024);
+			} else {
+				ac_declare_lds_as_pointer(&ctx->ac);
+				ctx->esgs_ring = ctx->ac.lds;
+			}
+		}
 	}
 
 	if (ctx->shader->is_gs_copy_shader) {
@@ -5129,7 +5012,7 @@ static void preload_ring_buffers(struct si_shader_context *ctx)
 
 			stride = 4 * num_components * sel->gs_max_out_vertices;
 
-			/* Limit on the stride field for <= CIK. */
+			/* Limit on the stride field for <= GFX7. */
 			assert(stride < (1 << 14));
 
 			num_records = 64;
@@ -5201,261 +5084,209 @@ static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx,
 	ac_build_kill_if_false(&ctx->ac, bit);
 }
 
-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 =
-		ac_shader_binary_config_start(binary, symbol_offset);
-	bool really_needs_scratch = false;
+/* For the UMR disassembler. */
+#define DEBUGGER_END_OF_CODE_MARKER	0xbf9f0000 /* invalid instruction */
+#define DEBUGGER_NUM_MARKERS		5
 
-	/* 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 ac_shader_reloc *reloc = &binary->relocs[i];
+static bool si_shader_binary_open(struct si_screen *screen,
+				  struct si_shader *shader,
+				  struct ac_rtld_binary *rtld)
+{
+	const struct si_shader_selector *sel = shader->selector;
+	enum pipe_shader_type shader_type = sel ? sel->type : PIPE_SHADER_COMPUTE;
+	const char *part_elfs[5];
+	size_t part_sizes[5];
+	unsigned num_parts = 0;
 
-		if (!strcmp(scratch_rsrc_dword0_symbol, reloc->name) ||
-		    !strcmp(scratch_rsrc_dword1_symbol, reloc->name)) {
-			really_needs_scratch = true;
-			break;
-		}
+#define add_part(shader_or_part) \
+	if (shader_or_part) { \
+		part_elfs[num_parts] = (shader_or_part)->binary.elf_buffer; \
+		part_sizes[num_parts] = (shader_or_part)->binary.elf_size; \
+		num_parts++; \
 	}
 
-	/* XXX: We may be able to emit some of these values directly rather than
-	 * extracting fields to be emitted later.
-	 */
+	add_part(shader->prolog);
+	add_part(shader->previous_stage);
+	add_part(shader->prolog2);
+	add_part(shader);
+	add_part(shader->epilog);
 
-	for (i = 0; i < binary->config_size_per_symbol; i+= 8) {
-		unsigned reg = util_le32_to_cpu(*(uint32_t*)(config + i));
-		unsigned value = util_le32_to_cpu(*(uint32_t*)(config + i + 4));
-		switch (reg) {
-		case R_00B028_SPI_SHADER_PGM_RSRC1_PS:
-		case R_00B128_SPI_SHADER_PGM_RSRC1_VS:
-		case R_00B228_SPI_SHADER_PGM_RSRC1_GS:
-		case R_00B428_SPI_SHADER_PGM_RSRC1_HS:
-		case R_00B848_COMPUTE_PGM_RSRC1:
-			conf->num_sgprs = MAX2(conf->num_sgprs, (G_00B028_SGPRS(value) + 1) * 8);
-			conf->num_vgprs = MAX2(conf->num_vgprs, (G_00B028_VGPRS(value) + 1) * 4);
-			conf->float_mode =  G_00B028_FLOAT_MODE(value);
-			conf->rsrc1 = value;
-			break;
-		case R_00B02C_SPI_SHADER_PGM_RSRC2_PS:
-			conf->lds_size = MAX2(conf->lds_size, G_00B02C_EXTRA_LDS_SIZE(value));
-			break;
-		case R_00B84C_COMPUTE_PGM_RSRC2:
-			conf->lds_size = MAX2(conf->lds_size, G_00B84C_LDS_SIZE(value));
-			conf->rsrc2 = value;
-			break;
-		case R_0286CC_SPI_PS_INPUT_ENA:
-			conf->spi_ps_input_ena = value;
-			break;
-		case R_0286D0_SPI_PS_INPUT_ADDR:
-			conf->spi_ps_input_addr = value;
-			break;
-		case R_0286E8_SPI_TMPRING_SIZE:
-		case R_00B860_COMPUTE_TMPRING_SIZE:
-			/* WAVESIZE is in units of 256 dwords. */
-			if (really_needs_scratch)
-				conf->scratch_bytes_per_wave =
-					G_00B860_WAVESIZE(value) * 256 * 4;
-			break;
-		case 0x4: /* SPILLED_SGPRS */
-			conf->spilled_sgprs = value;
-			break;
-		case 0x8: /* SPILLED_VGPRS */
-			conf->spilled_vgprs = value;
-			break;
-		default:
-			{
-				static bool printed;
+#undef add_part
 
-				if (!printed) {
-					fprintf(stderr, "Warning: LLVM emitted unknown "
-						"config register: 0x%x\n", reg);
-					printed = true;
-				}
-			}
-			break;
-		}
+	struct ac_rtld_symbol lds_symbols[1];
+	unsigned num_lds_symbols = 0;
+
+	if (sel && screen->info.chip_class >= GFX9 &&
+	    sel->type == PIPE_SHADER_GEOMETRY && !shader->is_gs_copy_shader) {
+		/* We add this symbol even on LLVM <= 8 to ensure that
+		 * shader->config.lds_size is set correctly below.
+		 */
+		struct ac_rtld_symbol *sym = &lds_symbols[num_lds_symbols++];
+		sym->name = "esgs_ring";
+		sym->size = shader->gs_info.esgs_ring_size;
+		sym->align = 64 * 1024;
+	}
+
+	bool ok = ac_rtld_open(rtld, (struct ac_rtld_open_info){
+			.info = &screen->info,
+			.options = {
+				.halt_at_entry = screen->options.halt_shaders,
+			},
+			.shader_type = tgsi_processor_to_shader_stage(shader_type),
+			.num_parts = num_parts,
+			.elf_ptrs = part_elfs,
+			.elf_sizes = part_sizes,
+			.num_shared_lds_symbols = num_lds_symbols,
+			.shared_lds_symbols = lds_symbols });
+
+	if (rtld->lds_size > 0) {
+		unsigned alloc_granularity = screen->info.chip_class >= GFX7 ? 512 : 256;
+		shader->config.lds_size =
+			align(rtld->lds_size, alloc_granularity) / alloc_granularity;
 	}
 
-	if (!conf->spi_ps_input_addr)
-		conf->spi_ps_input_addr = conf->spi_ps_input_ena;
+	return ok;
 }
 
-void si_shader_apply_scratch_relocs(struct si_shader *shader,
-				    uint64_t scratch_va)
+static unsigned si_get_shader_binary_size(struct si_screen *screen, struct si_shader *shader)
 {
-	unsigned i;
-	uint32_t scratch_rsrc_dword0 = scratch_va;
-	uint32_t scratch_rsrc_dword1 =
-		S_008F04_BASE_ADDRESS_HI(scratch_va >> 32);
-
-	/* Enable scratch coalescing. */
-	scratch_rsrc_dword1 |= S_008F04_SWIZZLE_ENABLE(1);
-
-	for (i = 0 ; i < shader->binary.reloc_count; i++) {
-		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,
-			&scratch_rsrc_dword0, 4);
-		} else if (!strcmp(scratch_rsrc_dword1_symbol, reloc->name)) {
-			util_memcpy_cpu_to_le32(shader->binary.code + reloc->offset,
-			&scratch_rsrc_dword1, 4);
-		}
-	}
+	struct ac_rtld_binary rtld;
+	si_shader_binary_open(screen, shader, &rtld);
+	return rtld.rx_size;
 }
 
-static unsigned si_get_shader_binary_size(const struct si_shader *shader)
+
+static bool si_get_external_symbol(void *data, const char *name, uint64_t *value)
 {
-	unsigned size = shader->binary.code_size;
+	uint64_t *scratch_va = data;
 
-	if (shader->prolog)
-		size += shader->prolog->binary.code_size;
-	if (shader->previous_stage)
-		size += shader->previous_stage->binary.code_size;
-	if (shader->prolog2)
-		size += shader->prolog2->binary.code_size;
-	if (shader->epilog)
-		size += shader->epilog->binary.code_size;
-	return size;
+	if (!strcmp(scratch_rsrc_dword0_symbol, name)) {
+		*value = (uint32_t)*scratch_va;
+		return true;
+	}
+	if (!strcmp(scratch_rsrc_dword1_symbol, name)) {
+		/* Enable scratch coalescing. */
+		*value = S_008F04_BASE_ADDRESS_HI(*scratch_va >> 32) |
+			 S_008F04_SWIZZLE_ENABLE(1);
+		if (HAVE_LLVM < 0x0800) {
+			/* Old LLVM created an R_ABS32_HI relocation for
+			 * this symbol. */
+			*value <<= 32;
+		}
+		return true;
+	}
+
+	return false;
 }
 
-int si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader)
+bool si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader,
+			     uint64_t scratch_va)
 {
-	const struct ac_shader_binary *prolog =
-		shader->prolog ? &shader->prolog->binary : NULL;
-	const struct ac_shader_binary *previous_stage =
-		shader->previous_stage ? &shader->previous_stage->binary : NULL;
-	const struct ac_shader_binary *prolog2 =
-		shader->prolog2 ? &shader->prolog2->binary : NULL;
-	const struct ac_shader_binary *epilog =
-		shader->epilog ? &shader->epilog->binary : NULL;
-	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;
-
-	assert(!prolog || !prolog->rodata_size);
-	assert(!previous_stage || !previous_stage->rodata_size);
-	assert(!prolog2 || !prolog2->rodata_size);
-	assert((!prolog && !previous_stage && !prolog2 && !epilog) ||
-	       !mainb->rodata_size);
-	assert(!epilog || !epilog->rodata_size);
+	struct ac_rtld_binary binary;
+	if (!si_shader_binary_open(sscreen, shader, &binary))
+		return false;
 
-	r600_resource_reference(&shader->bo, NULL);
-	shader->bo = (struct r600_resource*)
-		     si_aligned_buffer_create(&sscreen->b,
+	si_resource_reference(&shader->bo, NULL);
+	shader->bo = si_aligned_buffer_create(&sscreen->b,
 					      sscreen->cpdma_prefetch_writes_memory ?
 						0 : SI_RESOURCE_FLAG_READ_ONLY,
                                               PIPE_USAGE_IMMUTABLE,
-                                              align(bo_size, SI_CPDMA_ALIGNMENT),
+                                              align(binary.rx_size, SI_CPDMA_ALIGNMENT),
                                               256);
 	if (!shader->bo)
-		return -ENOMEM;
+		return false;
 
 	/* Upload. */
-	ptr = sscreen->ws->buffer_map(shader->bo->buf, NULL,
+	struct ac_rtld_upload_info u = {};
+	u.binary = &binary;
+	u.get_external_symbol = si_get_external_symbol;
+	u.cb_data = &scratch_va;
+	u.rx_va = shader->bo->gpu_address;
+	u.rx_ptr = sscreen->ws->buffer_map(shader->bo->buf, NULL,
 					PIPE_TRANSFER_READ_WRITE |
-					PIPE_TRANSFER_UNSYNCHRONIZED);
-
-	/* Don't use util_memcpy_cpu_to_le32. LLVM binaries are
-	 * endian-independent. */
-	if (prolog) {
-		memcpy(ptr, prolog->code, prolog->code_size);
-		ptr += prolog->code_size;
-	}
-	if (previous_stage) {
-		memcpy(ptr, previous_stage->code, previous_stage->code_size);
-		ptr += previous_stage->code_size;
-	}
-	if (prolog2) {
-		memcpy(ptr, prolog2->code, prolog2->code_size);
-		ptr += prolog2->code_size;
-	}
-
-	memcpy(ptr, mainb->code, mainb->code_size);
-	ptr += mainb->code_size;
+					PIPE_TRANSFER_UNSYNCHRONIZED |
+					RADEON_TRANSFER_TEMPORARY);
+	if (!u.rx_ptr)
+		return false;
 
-	if (epilog)
-		memcpy(ptr, epilog->code, epilog->code_size);
-	else if (mainb->rodata_size > 0)
-		memcpy(ptr, mainb->rodata, mainb->rodata_size);
+	bool ok = ac_rtld_upload(&u);
 
 	sscreen->ws->buffer_unmap(shader->bo->buf);
-	return 0;
+	ac_rtld_close(&binary);
+
+	return ok;
 }
 
-static void si_shader_dump_disassembly(const struct ac_shader_binary *binary,
+static void si_shader_dump_disassembly(struct si_screen *screen,
+				       const struct si_shader_binary *binary,
 				       struct pipe_debug_callback *debug,
 				       const char *name, FILE *file)
 {
-	char *line, *p;
-	unsigned i, count;
+	struct ac_rtld_binary rtld_binary;
 
-	if (binary->disasm_string) {
-		fprintf(file, "Shader %s disassembly:\n", name);
-		fprintf(file, "%s", binary->disasm_string);
+	if (!ac_rtld_open(&rtld_binary, (struct ac_rtld_open_info){
+			.info = &screen->info,
+			.num_parts = 1,
+			.elf_ptrs = &binary->elf_buffer,
+			.elf_sizes = &binary->elf_size }))
+		return;
 
-		if (debug && debug->debug_message) {
-			/* Very long debug messages are cut off, so send the
-			 * disassembly one line at a time. This causes more
-			 * overhead, but on the plus side it simplifies
-			 * parsing of resulting logs.
-			 */
-			pipe_debug_message(debug, SHADER_INFO,
-					   "Shader Disassembly Begin");
+	const char *disasm;
+	size_t nbytes;
 
-			line = binary->disasm_string;
-			while (*line) {
-				p = util_strchrnul(line, '\n');
-				count = p - line;
+	if (!ac_rtld_get_section_by_name(&rtld_binary, ".AMDGPU.disasm", &disasm, &nbytes))
+		goto out;
 
-				if (count) {
-					pipe_debug_message(debug, SHADER_INFO,
-							   "%.*s", count, line);
-				}
+	if (nbytes > INT_MAX)
+		goto out;
 
-				if (!*p)
-					break;
-				line = p + 1;
+	if (debug && debug->debug_message) {
+		/* Very long debug messages are cut off, so send the
+		 * disassembly one line at a time. This causes more
+		 * overhead, but on the plus side it simplifies
+		 * parsing of resulting logs.
+		 */
+		pipe_debug_message(debug, SHADER_INFO,
+				   "Shader Disassembly Begin");
+
+		uint64_t line = 0;
+		while (line < nbytes) {
+			int count = nbytes - line;
+			const char *nl = memchr(disasm + line, '\n', nbytes - line);
+			if (nl)
+				count = nl - (disasm + line);
+
+			if (count) {
+				pipe_debug_message(debug, SHADER_INFO,
+						   "%.*s", count, disasm + line);
 			}
 
-			pipe_debug_message(debug, SHADER_INFO,
-					   "Shader Disassembly End");
-		}
-	} else {
-		fprintf(file, "Shader %s binary:\n", name);
-		for (i = 0; i < binary->code_size; i += 4) {
-			fprintf(file, "@0x%x: %02x%02x%02x%02x\n", i,
-				binary->code[i + 3], binary->code[i + 2],
-				binary->code[i + 1], binary->code[i]);
+			line += count + 1;
 		}
+
+		pipe_debug_message(debug, SHADER_INFO,
+				   "Shader Disassembly End");
+	}
+
+	if (file) {
+		fprintf(file, "Shader %s disassembly:\n", name);
+		fprintf(file, "%*s", (int)nbytes, disasm);
 	}
+
+out:
+	ac_rtld_close(&rtld_binary);
 }
 
 static void si_calculate_max_simd_waves(struct si_shader *shader)
 {
 	struct si_screen *sscreen = shader->selector->screen;
-	struct si_shader_config *conf = &shader->config;
+	struct ac_shader_config *conf = &shader->config;
 	unsigned num_inputs = shader->selector->info.num_inputs;
-	unsigned lds_increment = sscreen->info.chip_class >= CIK ? 512 : 256;
+	unsigned lds_increment = sscreen->info.chip_class >= GFX7 ? 512 : 256;
 	unsigned lds_per_wave = 0;
 	unsigned max_simd_waves;
 
-	switch (sscreen->info.family) {
-	/* These always have 8 waves: */
-	case CHIP_POLARIS10:
-	case CHIP_POLARIS11:
-	case CHIP_POLARIS12:
-		max_simd_waves = 8;
-		break;
-	default:
-		max_simd_waves = 10;
-	}
+	max_simd_waves = ac_get_max_simd_waves(sscreen->info.family);
 
 	/* Compute LDS usage for PS. */
 	switch (shader->selector->type) {
@@ -5485,10 +5316,9 @@ static void si_calculate_max_simd_waves(struct si_shader *shader)
 
 	/* Compute the per-SIMD wave counts. */
 	if (conf->num_sgprs) {
-		if (sscreen->info.chip_class >= VI)
-			max_simd_waves = MIN2(max_simd_waves, 800 / conf->num_sgprs);
-		else
-			max_simd_waves = MIN2(max_simd_waves, 512 / conf->num_sgprs);
+		max_simd_waves =
+			MIN2(max_simd_waves,
+			     ac_get_num_physical_sgprs(sscreen->info.chip_class) / conf->num_sgprs);
 	}
 
 	if (conf->num_vgprs)
@@ -5499,32 +5329,36 @@ static void si_calculate_max_simd_waves(struct si_shader *shader)
 	if (lds_per_wave)
 		max_simd_waves = MIN2(max_simd_waves, 16384 / lds_per_wave);
 
-	conf->max_simd_waves = max_simd_waves;
+	shader->info.max_simd_waves = max_simd_waves;
 }
 
-void si_shader_dump_stats_for_shader_db(const struct si_shader *shader,
+void si_shader_dump_stats_for_shader_db(struct si_screen *screen,
+					struct si_shader *shader,
 					struct pipe_debug_callback *debug)
 {
-	const struct si_shader_config *conf = &shader->config;
+	const struct ac_shader_config *conf = &shader->config;
+
+	if (screen->options.debug_disassembly)
+		si_shader_dump_disassembly(screen, &shader->binary, debug, "main", NULL);
 
 	pipe_debug_message(debug, SHADER_INFO,
 			   "Shader Stats: SGPRS: %d VGPRS: %d Code Size: %d "
 			   "LDS: %d Scratch: %d Max Waves: %d Spilled SGPRs: %d "
 			   "Spilled VGPRs: %d PrivMem VGPRs: %d",
 			   conf->num_sgprs, conf->num_vgprs,
-			   si_get_shader_binary_size(shader),
+			   si_get_shader_binary_size(screen, shader),
 			   conf->lds_size, conf->scratch_bytes_per_wave,
-			   conf->max_simd_waves, conf->spilled_sgprs,
-			   conf->spilled_vgprs, conf->private_mem_vgprs);
+			   shader->info.max_simd_waves, conf->spilled_sgprs,
+			   conf->spilled_vgprs, shader->info.private_mem_vgprs);
 }
 
 static void si_shader_dump_stats(struct si_screen *sscreen,
-				 const struct si_shader *shader,
+				 struct si_shader *shader,
 			         unsigned processor,
 				 FILE *file,
 				 bool check_debug_option)
 {
-	const struct si_shader_config *conf = &shader->config;
+	const struct ac_shader_config *conf = &shader->config;
 
 	if (!check_debug_option ||
 	    si_can_dump_shader(sscreen, processor)) {
@@ -5548,10 +5382,10 @@ static void si_shader_dump_stats(struct si_screen *sscreen,
 			"********************\n\n\n",
 			conf->num_sgprs, conf->num_vgprs,
 			conf->spilled_sgprs, conf->spilled_vgprs,
-			conf->private_mem_vgprs,
-			si_get_shader_binary_size(shader),
+			shader->info.private_mem_vgprs,
+			si_get_shader_binary_size(sscreen, shader),
 			conf->lds_size, conf->scratch_bytes_per_wave,
-			conf->max_simd_waves);
+			shader->info.max_simd_waves);
 	}
 }
 
@@ -5563,6 +5397,8 @@ const char *si_get_shader_name(const struct si_shader *shader, unsigned processo
 			return "Vertex Shader as ES";
 		else if (shader->key.as_ls)
 			return "Vertex Shader as LS";
+		else if (shader->key.opt.vs_as_prim_discard_cs)
+			return "Vertex Shader as Primitive Discard CS";
 		else
 			return "Vertex Shader as VS";
 	case PIPE_SHADER_TESS_CTRL:
@@ -5586,7 +5422,7 @@ const char *si_get_shader_name(const struct si_shader *shader, unsigned processo
 	}
 }
 
-void si_shader_dump(struct si_screen *sscreen, const struct si_shader *shader,
+void si_shader_dump(struct si_screen *sscreen, struct si_shader *shader,
 		    struct pipe_debug_callback *debug, unsigned processor,
 		    FILE *file, bool check_debug_option)
 {
@@ -5613,19 +5449,19 @@ void si_shader_dump(struct si_screen *sscreen, const struct si_shader *shader,
 		fprintf(file, "\n%s:\n", si_get_shader_name(shader, processor));
 
 		if (shader->prolog)
-			si_shader_dump_disassembly(&shader->prolog->binary,
+			si_shader_dump_disassembly(sscreen, &shader->prolog->binary,
 						   debug, "prolog", file);
 		if (shader->previous_stage)
-			si_shader_dump_disassembly(&shader->previous_stage->binary,
+			si_shader_dump_disassembly(sscreen, &shader->previous_stage->binary,
 						   debug, "previous stage", file);
 		if (shader->prolog2)
-			si_shader_dump_disassembly(&shader->prolog2->binary,
+			si_shader_dump_disassembly(sscreen, &shader->prolog2->binary,
 						   debug, "prolog2", file);
 
-		si_shader_dump_disassembly(&shader->binary, debug, "main", file);
+		si_shader_dump_disassembly(sscreen, &shader->binary, debug, "main", file);
 
 		if (shader->epilog)
-			si_shader_dump_disassembly(&shader->epilog->binary,
+			si_shader_dump_disassembly(sscreen, &shader->epilog->binary,
 						   debug, "epilog", file);
 		fprintf(file, "\n");
 	}
@@ -5635,15 +5471,15 @@ void si_shader_dump(struct si_screen *sscreen, const struct si_shader *shader,
 }
 
 static int si_compile_llvm(struct si_screen *sscreen,
-			   struct ac_shader_binary *binary,
-			   struct si_shader_config *conf,
-			   LLVMTargetMachineRef tm,
+			   struct si_shader_binary *binary,
+			   struct ac_shader_config *conf,
+			   struct ac_llvm_compiler *compiler,
 			   LLVMModuleRef mod,
 			   struct pipe_debug_callback *debug,
 			   unsigned processor,
-			   const char *name)
+			   const char *name,
+			   bool less_optimized)
 {
-	int r = 0;
 	unsigned count = p_atomic_inc_return(&sscreen->num_compilations);
 
 	if (si_can_dump_shader(sscreen, processor)) {
@@ -5663,12 +5499,24 @@ static int si_compile_llvm(struct si_screen *sscreen,
 	}
 
 	if (!si_replace_shader(count, binary)) {
-		r = si_llvm_compile(mod, binary, tm, debug);
+		unsigned r = si_llvm_compile(mod, binary, compiler, debug,
+					     less_optimized);
 		if (r)
 			return r;
 	}
 
-	si_shader_binary_read_config(binary, conf, 0);
+	struct ac_rtld_binary rtld;
+	if (!ac_rtld_open(&rtld, (struct ac_rtld_open_info){
+			.info = &sscreen->info,
+			.num_parts = 1,
+			.elf_ptrs = &binary->elf_buffer,
+			.elf_sizes = &binary->elf_size }))
+		return -1;
+
+	bool ok = ac_rtld_read_config(&rtld, conf);
+	ac_rtld_close(&rtld);
+	if (!ok)
+		return -1;
 
 	/* Enable 64-bit and 16-bit denormals, because there is no performance
 	 * cost.
@@ -5680,28 +5528,11 @@ static int si_compile_llvm(struct si_screen *sscreen,
 	 * - Floating-point output modifiers would be ignored by the hw.
 	 * - Some opcodes don't support denormals, such as v_mad_f32. We would
 	 *   have to stop using those.
-	 * - SI & CI would be very slow.
+	 * - GFX6 & GFX7 would be very slow.
 	 */
 	conf->float_mode |= V_00B028_FP_64_DENORMS;
 
-	FREE(binary->config);
-	FREE(binary->global_symbol_offsets);
-	binary->config = NULL;
-	binary->global_symbol_offsets = NULL;
-
-	/* Some shaders can't have rodata because their binaries can be
-	 * concatenated.
-	 */
-	if (binary->rodata_size &&
-	    (processor == PIPE_SHADER_VERTEX ||
-	     processor == PIPE_SHADER_TESS_CTRL ||
-	     processor == PIPE_SHADER_TESS_EVAL ||
-	     processor == PIPE_SHADER_FRAGMENT)) {
-		fprintf(stderr, "radeonsi: The shader can't have rodata.");
-		return -EINVAL;
-	}
-
-	return r;
+	return 0;
 }
 
 static void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret)
@@ -5715,29 +5546,21 @@ static void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret)
 /* Generate code for the hardware VS shader stage to go with a geometry shader */
 struct si_shader *
 si_generate_gs_copy_shader(struct si_screen *sscreen,
-			   LLVMTargetMachineRef tm,
+			   struct ac_llvm_compiler *compiler,
 			   struct si_shader_selector *gs_selector,
 			   struct pipe_debug_callback *debug)
 {
 	struct si_shader_context ctx;
 	struct si_shader *shader;
 	LLVMBuilderRef builder;
-	struct lp_build_tgsi_context *bld_base = &ctx.bld_base;
-	struct lp_build_context *uint = &bld_base->uint_bld;
-	struct si_shader_output_values *outputs;
+	struct si_shader_output_values outputs[SI_MAX_VS_OUTPUTS];
 	struct tgsi_shader_info *gsinfo = &gs_selector->info;
-	int i, r;
-
-	outputs = MALLOC(gsinfo->num_outputs * sizeof(outputs[0]));
+	int i;
 
-	if (!outputs)
-		return NULL;
 
 	shader = CALLOC_STRUCT(si_shader);
-	if (!shader) {
-		FREE(outputs);
+	if (!shader)
 		return NULL;
-	}
 
 	/* We can leave the fence as permanently signaled because the GS copy
 	 * shader only becomes visible globally after it has been compiled. */
@@ -5746,7 +5569,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
 	shader->selector = gs_selector;
 	shader->is_gs_copy_shader = true;
 
-	si_init_shader_ctx(&ctx, sscreen, tm);
+	si_init_shader_ctx(&ctx, sscreen, compiler);
 	ctx.shader = shader;
 	ctx.type = PIPE_SHADER_VERTEX;
 
@@ -5756,7 +5579,8 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
 	preload_ring_buffers(&ctx);
 
 	LLVMValueRef voffset =
-		lp_build_mul_imm(uint, ctx.abi.vertex_id, 4);
+		LLVMBuildMul(ctx.ac.builder, ctx.abi.vertex_id,
+			     LLVMConstInt(ctx.i32, 4, 0), "");
 
 	/* Fetch the vertex stream ID.*/
 	LLVMValueRef stream_id;
@@ -5803,7 +5627,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
 			for (unsigned chan = 0; chan < 4; chan++) {
 				if (!(gsinfo->output_usagemask[i] & (1 << chan)) ||
 				    outputs[i].vertex_stream[chan] != stream) {
-					outputs[i].values[chan] = ctx.bld_base.base.undef;
+					outputs[i].values[chan] = LLVMGetUndef(ctx.f32);
 					continue;
 				}
 
@@ -5840,26 +5664,30 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
 	ctx.type = PIPE_SHADER_GEOMETRY; /* override for shader dumping */
 	si_llvm_optimize_module(&ctx);
 
-	r = si_compile_llvm(sscreen, &ctx.shader->binary,
-			    &ctx.shader->config, ctx.tm,
-			    ctx.gallivm.module,
+	bool ok = false;
+	if (si_compile_llvm(sscreen, &ctx.shader->binary,
+			    &ctx.shader->config, ctx.compiler,
+			    ctx.ac.module,
 			    debug, PIPE_SHADER_GEOMETRY,
-			    "GS Copy Shader");
-	if (!r) {
+			    "GS Copy Shader", false) == 0) {
 		if (si_can_dump_shader(sscreen, PIPE_SHADER_GEOMETRY))
 			fprintf(stderr, "GS Copy Shader:\n");
 		si_shader_dump(sscreen, ctx.shader, debug,
 			       PIPE_SHADER_GEOMETRY, stderr, true);
-		r = si_shader_binary_upload(sscreen, ctx.shader);
+
+		if (!ctx.shader->config.scratch_bytes_per_wave)
+			ok = si_shader_binary_upload(sscreen, ctx.shader, 0);
+		else
+			ok = true;
 	}
 
 	si_llvm_dispose(&ctx);
 
-	FREE(outputs);
-
-	if (r != 0) {
+	if (!ok) {
 		FREE(shader);
 		shader = NULL;
+	} else {
+		si_fix_resource_usage(sscreen, shader);
 	}
 	return shader;
 }
@@ -5872,12 +5700,23 @@ static void si_dump_shader_key_vs(const struct si_shader_key *key,
 		prefix, prolog->instance_divisor_is_one);
 	fprintf(f, "  %s.instance_divisor_is_fetched = %u\n",
 		prefix, prolog->instance_divisor_is_fetched);
+	fprintf(f, "  %s.unpack_instance_id_from_vertex_id = %u\n",
+		prefix, prolog->unpack_instance_id_from_vertex_id);
 	fprintf(f, "  %s.ls_vgpr_fix = %u\n",
 		prefix, prolog->ls_vgpr_fix);
 
+	fprintf(f, "  mono.vs.fetch_opencode = %x\n", key->mono.vs_fetch_opencode);
 	fprintf(f, "  mono.vs.fix_fetch = {");
-	for (int i = 0; i < SI_MAX_ATTRIBS; i++)
-		fprintf(f, !i ? "%u" : ", %u", key->mono.vs_fix_fetch[i]);
+	for (int i = 0; i < SI_MAX_ATTRIBS; i++) {
+		union si_vs_fix_fetch fix = key->mono.vs_fix_fetch[i];
+		if (i)
+			fprintf(f, ", ");
+		if (!fix.bits)
+			fprintf(f, "0");
+		else
+			fprintf(f, "%u.%u.%u.%u", fix.u.reverse, fix.u.log_size,
+				fix.u.num_channels_m1, fix.u.format);
+	}
 	fprintf(f, "}\n");
 }
 
@@ -5896,6 +5735,28 @@ static void si_dump_shader_key(unsigned processor, const struct si_shader *shade
 		fprintf(f, "  as_ls = %u\n", key->as_ls);
 		fprintf(f, "  mono.u.vs_export_prim_id = %u\n",
 			key->mono.u.vs_export_prim_id);
+		fprintf(f, "  opt.vs_as_prim_discard_cs = %u\n",
+			key->opt.vs_as_prim_discard_cs);
+		fprintf(f, "  opt.cs_prim_type = %s\n",
+			tgsi_primitive_names[key->opt.cs_prim_type]);
+		fprintf(f, "  opt.cs_indexed = %u\n",
+			key->opt.cs_indexed);
+		fprintf(f, "  opt.cs_instancing = %u\n",
+			key->opt.cs_instancing);
+		fprintf(f, "  opt.cs_primitive_restart = %u\n",
+			key->opt.cs_primitive_restart);
+		fprintf(f, "  opt.cs_provoking_vertex_first = %u\n",
+			key->opt.cs_provoking_vertex_first);
+		fprintf(f, "  opt.cs_need_correct_orientation = %u\n",
+			key->opt.cs_need_correct_orientation);
+		fprintf(f, "  opt.cs_cull_front = %u\n",
+			key->opt.cs_cull_front);
+		fprintf(f, "  opt.cs_cull_back = %u\n",
+			key->opt.cs_cull_back);
+		fprintf(f, "  opt.cs_cull_z = %u\n",
+			key->opt.cs_cull_z);
+		fprintf(f, "  opt.cs_halfz_clip_space = %u\n",
+			key->opt.cs_halfz_clip_space);
 		break;
 
 	case PIPE_SHADER_TESS_CTRL:
@@ -5963,18 +5824,18 @@ static void si_dump_shader_key(unsigned processor, const struct si_shader *shade
 
 static void si_init_shader_ctx(struct si_shader_context *ctx,
 			       struct si_screen *sscreen,
-			       LLVMTargetMachineRef tm)
+			       struct ac_llvm_compiler *compiler)
 {
 	struct lp_build_tgsi_context *bld_base;
 
-	si_llvm_context_init(ctx, sscreen, tm);
+	si_llvm_context_init(ctx, sscreen, compiler);
 
 	bld_base = &ctx->bld_base;
 	bld_base->emit_fetch_funcs[TGSI_FILE_CONSTANT] = fetch_constant;
 
-	bld_base->op_actions[TGSI_OPCODE_INTERP_CENTROID] = interp_action;
-	bld_base->op_actions[TGSI_OPCODE_INTERP_SAMPLE] = interp_action;
-	bld_base->op_actions[TGSI_OPCODE_INTERP_OFFSET] = interp_action;
+	bld_base->op_actions[TGSI_OPCODE_INTERP_CENTROID].emit = build_interp_intrinsic;
+	bld_base->op_actions[TGSI_OPCODE_INTERP_SAMPLE].emit = build_interp_intrinsic;
+	bld_base->op_actions[TGSI_OPCODE_INTERP_OFFSET].emit = build_interp_intrinsic;
 
 	bld_base->op_actions[TGSI_OPCODE_MEMBAR].emit = membar_emit;
 
@@ -5992,7 +5853,6 @@ static void si_init_shader_ctx(struct si_shader_context *ctx,
 	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_tgsi_emit_vertex;
@@ -6025,9 +5885,9 @@ static void si_init_exec_from_input(struct si_shader_context *ctx,
 		LLVMGetParam(ctx->main_fn, param),
 		LLVMConstInt(ctx->i32, bitoffset, 0),
 	};
-	lp_build_intrinsic(ctx->ac.builder,
+	ac_build_intrinsic(&ctx->ac,
 			   "llvm.amdgcn.init.exec.from.input",
-			   ctx->voidt, args, 2, LP_FUNC_ATTR_CONVERGENT);
+			   ctx->voidt, args, 2, AC_FUNC_ATTR_CONVERGENT);
 }
 
 static bool si_vs_needs_prolog(const struct si_shader_selector *sel,
@@ -6038,8 +5898,7 @@ static bool si_vs_needs_prolog(const struct si_shader_selector *sel,
 	return sel->vs_needs_prolog || key->ls_vgpr_fix;
 }
 
-static bool si_compile_tgsi_main(struct si_shader_context *ctx,
-				 bool is_monolithic)
+static bool si_compile_tgsi_main(struct si_shader_context *ctx)
 {
 	struct si_shader *shader = ctx->shader;
 	struct si_shader_selector *sel = shader->selector;
@@ -6053,6 +5912,8 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx,
 			ctx->abi.emit_outputs = si_llvm_emit_ls_epilogue;
 		else if (shader->key.as_es)
 			ctx->abi.emit_outputs = si_llvm_emit_es_epilogue;
+		else if (shader->key.opt.vs_as_prim_discard_cs)
+			ctx->abi.emit_outputs = si_llvm_emit_prim_discard_cs_epilogue;
 		else
 			ctx->abi.emit_outputs = si_llvm_emit_vs_epilogue;
 		bld_base->emit_epilogue = si_tgsi_emit_epilogue;
@@ -6124,7 +5985,7 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx,
 	 * if-block together with its prolog in si_build_wrapper_function.
 	 */
 	if (ctx->screen->info.chip_class >= GFX9) {
-		if (!is_monolithic &&
+		if (!shader->is_monolithic &&
 		    sel->info.num_instructions > 1 && /* not empty shader */
 		    (shader->key.as_es || shader->key.as_ls) &&
 		    (ctx->type == PIPE_SHADER_TESS_EVAL ||
@@ -6134,19 +5995,27 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx,
 						ctx->param_merged_wave_info, 0);
 		} else if (ctx->type == PIPE_SHADER_TESS_CTRL ||
 			   ctx->type == PIPE_SHADER_GEOMETRY) {
-			if (!is_monolithic)
+			if (!shader->is_monolithic)
 				ac_init_exec_full_mask(&ctx->ac);
 
-			/* The barrier must execute for all shaders in a
-			 * threadgroup.
-			 */
-			si_llvm_emit_barrier(NULL, bld_base, NULL);
-
 			LLVMValueRef num_threads = si_unpack_param(ctx, ctx->param_merged_wave_info, 8, 8);
 			LLVMValueRef ena =
 				LLVMBuildICmp(ctx->ac.builder, LLVMIntULT,
 					    ac_get_thread_id(&ctx->ac), num_threads, "");
 			lp_build_if(&ctx->merged_wrap_if_state, &ctx->gallivm, ena);
+
+			/* The barrier must execute for all shaders in a
+			 * threadgroup.
+			 *
+			 * Execute the barrier inside the conditional block,
+			 * so that empty waves can jump directly to s_endpgm,
+			 * which will also signal the barrier.
+			 *
+			 * If the shader is TCS and the TCS epilog is present
+			 * and contains a barrier, it will wait there and then
+			 * reach s_endpgm.
+			 */
+			si_llvm_emit_barrier(NULL, bld_base, NULL);
 		}
 	}
 
@@ -6154,7 +6023,7 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx,
 	    sel->tcs_info.tessfactors_are_def_in_all_invocs) {
 		for (unsigned i = 0; i < 6; i++) {
 			ctx->invoc0_tess_factors[i] =
-				lp_build_alloca_undef(&ctx->gallivm, ctx->i32, "");
+				ac_build_alloca_undef(&ctx->ac, ctx->i32, "");
 		}
 	}
 
@@ -6162,15 +6031,14 @@ 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(&ctx->gallivm,
-						ctx->i32, "");
+				ac_build_alloca(&ctx->ac, ctx->i32, "");
 		}
 	}
 
 	if (sel->force_correct_derivs_after_kill) {
-		ctx->postponed_kill = lp_build_alloca_undef(&ctx->gallivm, ctx->i1, "");
+		ctx->postponed_kill = ac_build_alloca_undef(&ctx->ac, ctx->i1, "");
 		/* true = don't kill. */
-		LLVMBuildStore(ctx->ac.builder, LLVMConstInt(ctx->i1, 1, 0),
+		LLVMBuildStore(ctx->ac.builder, ctx->i1true,
 			       ctx->postponed_kill);
 	}
 
@@ -6261,7 +6129,8 @@ static void si_get_ps_prolog_key(struct si_shader *shader,
 			/* BCOLORs are stored after the last input. */
 			key->ps_prolog.num_interp_inputs = info->num_inputs;
 			key->ps_prolog.face_vgpr_index = shader->info.face_vgpr_index;
-			shader->config.spi_ps_input_ena |= S_0286CC_FRONT_FACE_ENA(1);
+			if (separate_prolog)
+				shader->config.spi_ps_input_ena |= S_0286CC_FRONT_FACE_ENA(1);
 		}
 
 		for (unsigned i = 0; i < 2; i++) {
@@ -6292,18 +6161,24 @@ static void si_get_ps_prolog_key(struct si_shader *shader,
 				switch (location) {
 				case TGSI_INTERPOLATE_LOC_SAMPLE:
 					key->ps_prolog.color_interp_vgpr_index[i] = 0;
-					shader->config.spi_ps_input_ena |=
-						S_0286CC_PERSP_SAMPLE_ENA(1);
+					if (separate_prolog) {
+						shader->config.spi_ps_input_ena |=
+							S_0286CC_PERSP_SAMPLE_ENA(1);
+					}
 					break;
 				case TGSI_INTERPOLATE_LOC_CENTER:
 					key->ps_prolog.color_interp_vgpr_index[i] = 2;
-					shader->config.spi_ps_input_ena |=
-						S_0286CC_PERSP_CENTER_ENA(1);
+					if (separate_prolog) {
+						shader->config.spi_ps_input_ena |=
+							S_0286CC_PERSP_CENTER_ENA(1);
+					}
 					break;
 				case TGSI_INTERPOLATE_LOC_CENTROID:
 					key->ps_prolog.color_interp_vgpr_index[i] = 4;
-					shader->config.spi_ps_input_ena |=
-						S_0286CC_PERSP_CENTROID_ENA(1);
+					if (separate_prolog) {
+						shader->config.spi_ps_input_ena |=
+							S_0286CC_PERSP_CENTROID_ENA(1);
+					}
 					break;
 				default:
 					assert(0);
@@ -6324,20 +6199,26 @@ static void si_get_ps_prolog_key(struct si_shader *shader,
 				case TGSI_INTERPOLATE_LOC_SAMPLE:
 					key->ps_prolog.color_interp_vgpr_index[i] =
 						separate_prolog ? 6 : 9;
-					shader->config.spi_ps_input_ena |=
-						S_0286CC_LINEAR_SAMPLE_ENA(1);
+					if (separate_prolog) {
+						shader->config.spi_ps_input_ena |=
+							S_0286CC_LINEAR_SAMPLE_ENA(1);
+					}
 					break;
 				case TGSI_INTERPOLATE_LOC_CENTER:
 					key->ps_prolog.color_interp_vgpr_index[i] =
 						separate_prolog ? 8 : 11;
-					shader->config.spi_ps_input_ena |=
-						S_0286CC_LINEAR_CENTER_ENA(1);
+					if (separate_prolog) {
+						shader->config.spi_ps_input_ena |=
+							S_0286CC_LINEAR_CENTER_ENA(1);
+					}
 					break;
 				case TGSI_INTERPOLATE_LOC_CENTROID:
 					key->ps_prolog.color_interp_vgpr_index[i] =
 						separate_prolog ? 10 : 13;
-					shader->config.spi_ps_input_ena |=
-						S_0286CC_LINEAR_CENTROID_ENA(1);
+					if (separate_prolog) {
+						shader->config.spi_ps_input_ena |=
+							S_0286CC_LINEAR_CENTROID_ENA(1);
+					}
 					break;
 				default:
 					assert(0);
@@ -6535,7 +6416,8 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
 	si_init_function_info(&fninfo);
 
 	for (unsigned i = 0; i < num_parts; ++i) {
-		lp_add_function_attr(parts[i], -1, LP_FUNC_ATTR_ALWAYSINLINE);
+		ac_add_function_attr(ctx->ac.context, parts[i], -1,
+				     AC_FUNC_ATTR_ALWAYSINLINE);
 		LLVMSetLinkage(parts[i], LLVMPrivateLinkage);
 	}
 
@@ -6576,10 +6458,29 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
 		gprs += size;
 	}
 
-	si_create_function(ctx, "wrapper", NULL, 0, &fninfo,
+	/* Prepare the return type. */
+	unsigned num_returns = 0;
+	LLVMTypeRef returns[32], last_func_type, return_type;
+
+	last_func_type = LLVMGetElementType(LLVMTypeOf(parts[num_parts - 1]));
+	return_type = LLVMGetReturnType(last_func_type);
+
+	switch (LLVMGetTypeKind(return_type)) {
+	case LLVMStructTypeKind:
+		num_returns = LLVMCountStructElementTypes(return_type);
+		assert(num_returns <= ARRAY_SIZE(returns));
+		LLVMGetStructElementTypes(return_type, returns);
+		break;
+	case LLVMVoidTypeKind:
+		break;
+	default:
+		unreachable("unexpected type");
+	}
+
+	si_create_function(ctx, "wrapper", returns, num_returns, &fninfo,
 			   si_get_max_workgroup_size(ctx->shader));
 
-	if (is_merged_shader(ctx->shader))
+	if (is_merged_shader(ctx))
 		ac_init_exec_full_mask(&ctx->ac);
 
 	/* Record the arguments of the function as if they were an output of
@@ -6628,16 +6529,16 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
 	initial_num_out_sgpr = num_out_sgpr;
 
 	/* Now chain the parts. */
+	LLVMValueRef ret;
 	for (unsigned part = 0; part < num_parts; ++part) {
 		LLVMValueRef in[48];
-		LLVMValueRef ret;
 		LLVMTypeRef ret_type;
 		unsigned out_idx = 0;
 		unsigned num_params = LLVMCountParams(parts[part]);
 
 		/* Merged shaders are executed conditionally depending
 		 * on the number of enabled threads passed in the input SGPRs. */
-		if (is_merged_shader(ctx->shader) && part == 0) {
+		if (is_merged_shader(ctx) && part == 0) {
 			LLVMValueRef ena, count = initial[3];
 
 			count = LLVMBuildAnd(builder, count,
@@ -6662,9 +6563,10 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
 			param_size = ac_get_type_size(param_type) / 4;
 			is_sgpr = ac_is_sgpr_param(param);
 
-			if (is_sgpr)
-				lp_add_function_attr(parts[part], param_idx + 1, LP_FUNC_ATTR_INREG);
-			else if (out_idx < num_out_sgpr) {
+			if (is_sgpr) {
+				ac_add_function_attr(ctx->ac.context, parts[part],
+						     param_idx + 1, AC_FUNC_ATTR_INREG);
+			} else if (out_idx < num_out_sgpr) {
 				/* Skip returned SGPRs the current part doesn't
 				 * declare on the input. */
 				out_idx = num_out_sgpr;
@@ -6675,12 +6577,12 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
 			if (param_size == 1)
 				arg = out[out_idx];
 			else
-				arg = lp_build_gather_values(&ctx->gallivm, &out[out_idx], param_size);
+				arg = ac_build_gather_values(&ctx->ac, &out[out_idx], param_size);
 
 			if (LLVMTypeOf(arg) != param_type) {
 				if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
 					if (LLVMGetPointerAddressSpace(param_type) ==
-					    AC_CONST_32BIT_ADDR_SPACE) {
+					    AC_ADDR_SPACE_CONST_32BIT) {
 						arg = LLVMBuildBitCast(builder, arg, ctx->i32, "");
 						arg = LLVMBuildIntToPtr(builder, arg, param_type, "");
 					} else {
@@ -6696,9 +6598,9 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
 			out_idx += param_size;
 		}
 
-		ret = LLVMBuildCall(builder, parts[part], in, num_params, "");
+		ret = ac_build_call(&ctx->ac, parts[part], in, num_params);
 
-		if (is_merged_shader(ctx->shader) &&
+		if (is_merged_shader(ctx) &&
 		    part + 1 == next_shader_first_part) {
 			lp_build_endif(&if_state);
 
@@ -6741,13 +6643,32 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
 		}
 	}
 
-	LLVMBuildRetVoid(builder);
+	/* Return the value from the last part. */
+	if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind)
+		LLVMBuildRetVoid(builder);
+	else
+		LLVMBuildRet(builder, ret);
+}
+
+static bool si_should_optimize_less(struct ac_llvm_compiler *compiler,
+				    struct si_shader_selector *sel)
+{
+	if (!compiler->low_opt_passes)
+		return false;
+
+	/* Assume a slow CPU. */
+	assert(!sel->screen->info.has_dedicated_vram &&
+	       sel->screen->info.chip_class <= GFX8);
+
+	/* For a crazy dEQP test containing 2597 memory opcodes, mostly
+	 * buffer stores. */
+	return sel->type == PIPE_SHADER_COMPUTE &&
+	       sel->info.num_memory_instructions > 1000;
 }
 
 int si_compile_tgsi_shader(struct si_screen *sscreen,
-			   LLVMTargetMachineRef tm,
+			   struct ac_llvm_compiler *compiler,
 			   struct si_shader *shader,
-			   bool is_monolithic,
 			   struct pipe_debug_callback *debug)
 {
 	struct si_shader_selector *sel = shader->selector;
@@ -6765,21 +6686,20 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 		si_dump_streamout(&sel->so);
 	}
 
-	si_init_shader_ctx(&ctx, sscreen, tm);
+	si_init_shader_ctx(&ctx, sscreen, compiler);
 	si_llvm_context_set_tgsi(&ctx, shader);
-	ctx.separate_prolog = !is_monolithic;
 
 	memset(shader->info.vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
 	       sizeof(shader->info.vs_output_param_offset));
 
 	shader->info.uses_instanceid = sel->info.uses_instanceid;
 
-	if (!si_compile_tgsi_main(&ctx, is_monolithic)) {
+	if (!si_compile_tgsi_main(&ctx)) {
 		si_llvm_dispose(&ctx);
 		return -1;
 	}
 
-	if (is_monolithic && ctx.type == PIPE_SHADER_VERTEX) {
+	if (shader->is_monolithic && ctx.type == PIPE_SHADER_VERTEX) {
 		LLVMValueRef parts[2];
 		bool need_prolog = sel->vs_needs_prolog;
 
@@ -6797,7 +6717,10 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 
 		si_build_wrapper_function(&ctx, parts + !need_prolog,
 					  1 + need_prolog, need_prolog, 0);
-	} else if (is_monolithic && ctx.type == PIPE_SHADER_TESS_CTRL) {
+
+		if (ctx.shader->key.opt.vs_as_prim_discard_cs)
+			si_build_prim_discard_compute_shader(&ctx);
+	} else if (shader->is_monolithic && ctx.type == PIPE_SHADER_TESS_CTRL) {
 		if (sscreen->info.chip_class >= GFX9) {
 			struct si_shader_selector *ls = shader->key.part.tcs.ls;
 			LLVMValueRef parts[4];
@@ -6820,9 +6743,10 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 			shader_ls.key.as_ls = 1;
 			shader_ls.key.mono = shader->key.mono;
 			shader_ls.key.opt = shader->key.opt;
+			shader_ls.is_monolithic = true;
 			si_llvm_context_set_tgsi(&ctx, &shader_ls);
 
-			if (!si_compile_tgsi_main(&ctx, true)) {
+			if (!si_compile_tgsi_main(&ctx)) {
 				si_llvm_dispose(&ctx);
 				return -1;
 			}
@@ -6862,7 +6786,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 
 			si_build_wrapper_function(&ctx, parts, 2, 0, 0);
 		}
-	} else if (is_monolithic && ctx.type == PIPE_SHADER_GEOMETRY) {
+	} else if (shader->is_monolithic && ctx.type == PIPE_SHADER_GEOMETRY) {
 		if (ctx.screen->info.chip_class >= GFX9) {
 			struct si_shader_selector *es = shader->key.part.gs.es;
 			LLVMValueRef es_prolog = NULL;
@@ -6884,9 +6808,10 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 			shader_es.key.as_es = 1;
 			shader_es.key.mono = shader->key.mono;
 			shader_es.key.opt = shader->key.opt;
+			shader_es.is_monolithic = true;
 			si_llvm_context_set_tgsi(&ctx, &shader_es);
 
-			if (!si_compile_tgsi_main(&ctx, true)) {
+			if (!si_compile_tgsi_main(&ctx)) {
 				si_llvm_dispose(&ctx);
 				return -1;
 			}
@@ -6935,7 +6860,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 
 			si_build_wrapper_function(&ctx, parts, 2, 1, 0);
 		}
-	} else if (is_monolithic && ctx.type == PIPE_SHADER_FRAGMENT) {
+	} else if (shader->is_monolithic && ctx.type == PIPE_SHADER_FRAGMENT) {
 		LLVMValueRef parts[3];
 		union si_shader_part_key prolog_key;
 		union si_shader_part_key epilog_key;
@@ -6966,7 +6891,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 
 	if ((debug && debug->debug_message) ||
 	    si_can_dump_shader(sscreen, ctx.type)) {
-		ctx.shader->config.private_mem_vgprs =
+		ctx.shader->info.private_mem_vgprs =
 			ac_count_scratch_private_memory(ctx.main_fn);
 	}
 
@@ -6975,8 +6900,10 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 	       LLVMPointerTypeKind);
 
 	/* Compile to bytecode. */
-	r = si_compile_llvm(sscreen, &shader->binary, &shader->config, tm,
-			    ctx.gallivm.module, debug, ctx.type, "TGSI shader");
+	r = si_compile_llvm(sscreen, &shader->binary, &shader->config, compiler,
+			    ctx.ac.module, debug, ctx.type,
+			    si_get_shader_name(shader, ctx.type),
+			    si_should_optimize_less(compiler, shader->selector));
 	si_llvm_dispose(&ctx);
 	if (r) {
 		fprintf(stderr, "LLVM failed to compile shader\n");
@@ -6989,7 +6916,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 	if (sel->type == PIPE_SHADER_COMPUTE) {
 		unsigned wave_size = 64;
 		unsigned max_vgprs = 256;
-		unsigned max_sgprs = sscreen->info.chip_class >= VI ? 800 : 512;
+		unsigned max_sgprs = sscreen->info.chip_class >= GFX8 ? 800 : 512;
 		unsigned max_sgprs_per_wave = 128;
 		unsigned max_block_threads = si_get_max_workgroup_size(shader);
 		unsigned min_waves_per_cu = DIV_ROUND_UP(max_block_threads, wave_size);
@@ -7015,7 +6942,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 	}
 
 	/* Add the scratch offset to input SGPRs. */
-	if (shader->config.scratch_bytes_per_wave && !is_merged_shader(shader))
+	if (shader->config.scratch_bytes_per_wave && !is_merged_shader(&ctx))
 		shader->info.num_input_sgprs += 1; /* scratch byte offset */
 
 	/* Calculate the number of fragment input VGPRs. */
@@ -7063,7 +6990,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
 	}
 
 	si_calculate_max_simd_waves(shader);
-	si_shader_dump_stats_for_shader_db(shader, debug);
+	si_shader_dump_stats_for_shader_db(sscreen, shader, debug);
 	return 0;
 }
 
@@ -7086,7 +7013,7 @@ si_get_shader_part(struct si_screen *sscreen,
 		   enum pipe_shader_type type,
 		   bool prolog,
 		   union si_shader_part_key *key,
-		   LLVMTargetMachineRef tm,
+		   struct ac_llvm_compiler *compiler,
 		   struct pipe_debug_callback *debug,
 		   void (*build)(struct si_shader_context *,
 				 union si_shader_part_key *),
@@ -7111,7 +7038,7 @@ si_get_shader_part(struct si_screen *sscreen,
 	struct si_shader shader = {};
 	struct si_shader_context ctx;
 
-	si_init_shader_ctx(&ctx, sscreen, tm);
+	si_init_shader_ctx(&ctx, sscreen, compiler);
 	ctx.shader = &shader;
 	ctx.type = type;
 
@@ -7142,8 +7069,8 @@ si_get_shader_part(struct si_screen *sscreen,
 	/* Compile. */
 	si_llvm_optimize_module(&ctx);
 
-	if (si_compile_llvm(sscreen, &result->binary, &result->config, tm,
-			    ctx.ac.module, debug, ctx.type, name)) {
+	if (si_compile_llvm(sscreen, &result->binary, &result->config, compiler,
+			    ctx.ac.module, debug, ctx.type, name, false)) {
 		FREE(result);
 		result = NULL;
 		goto out;
@@ -7161,26 +7088,11 @@ out:
 static LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx)
 {
 	LLVMValueRef ptr[2], list;
-	bool is_merged_shader =
-		ctx->screen->info.chip_class >= GFX9 &&
-		(ctx->type == PIPE_SHADER_TESS_CTRL ||
-		 ctx->type == PIPE_SHADER_GEOMETRY ||
-		 ctx->shader->key.as_ls || ctx->shader->key.as_es);
-
-	if (HAVE_32BIT_POINTERS) {
-		ptr[0] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
-		list = LLVMBuildIntToPtr(ctx->ac.builder, ptr[0],
-					 ac_array_in_const32_addr_space(ctx->v4i32), "");
-		return list;
-	}
-
-	/* Get the pointer to rw buffers. */
-	ptr[0] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
-	ptr[1] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS + 1);
-	list = lp_build_gather_values(&ctx->gallivm, ptr, 2);
-	list = LLVMBuildBitCast(ctx->ac.builder, list, ctx->i64, "");
-	list = LLVMBuildIntToPtr(ctx->ac.builder, list,
-				 ac_array_in_const_addr_space(ctx->v4i32), "");
+	bool merged_shader = is_merged_shader(ctx);
+
+	ptr[0] = LLVMGetParam(ctx->main_fn, (merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
+	list = LLVMBuildIntToPtr(ctx->ac.builder, ptr[0],
+				 ac_array_in_const32_addr_space(ctx->v4i32), "");
 	return list;
 }
 
@@ -7265,8 +7177,21 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx,
 		}
 	}
 
-	ctx->abi.vertex_id = input_vgprs[first_vs_vgpr];
-	ctx->abi.instance_id = input_vgprs[first_vs_vgpr + (key->vs_prolog.as_ls ? 2 : 1)];
+	unsigned vertex_id_vgpr = first_vs_vgpr;
+	unsigned instance_id_vgpr = first_vs_vgpr + (key->vs_prolog.as_ls ? 2 : 1);
+
+	ctx->abi.vertex_id = input_vgprs[vertex_id_vgpr];
+	ctx->abi.instance_id = input_vgprs[instance_id_vgpr];
+
+	/* InstanceID = VertexID >> 16;
+	 * VertexID   = VertexID & 0xffff;
+	 */
+	if (key->vs_prolog.states.unpack_instance_id_from_vertex_id) {
+		ctx->abi.instance_id = LLVMBuildLShr(ctx->ac.builder, ctx->abi.vertex_id,
+						     LLVMConstInt(ctx->i32, 16, 0), "");
+		ctx->abi.vertex_id = LLVMBuildAnd(ctx->ac.builder, ctx->abi.vertex_id,
+						  LLVMConstInt(ctx->i32, 0xffff, 0), "");
+	}
 
 	/* Copy inputs to outputs. This should be no-op, as the registers match,
 	 * but it will prevent the compiler from overwriting them unintentionally.
@@ -7278,6 +7203,12 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx,
 	}
 	for (i = 0; i < num_input_vgprs; i++) {
 		LLVMValueRef p = input_vgprs[i];
+
+		if (i == vertex_id_vgpr)
+			p = ctx->abi.vertex_id;
+		else if (i == instance_id_vgpr)
+			p = ctx->abi.instance_id;
+
 		p = ac_to_float(&ctx->ac, p);
 		ret = LLVMBuildInsertValue(ctx->ac.builder, ret, p,
 					   key->vs_prolog.num_input_sgprs + i, "");
@@ -7299,22 +7230,32 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx,
 			key->vs_prolog.states.instance_divisor_is_one & (1u << i);
 		bool divisor_is_fetched =
 			key->vs_prolog.states.instance_divisor_is_fetched & (1u << i);
-		LLVMValueRef index;
-
-		if (divisor_is_one || divisor_is_fetched) {
-			LLVMValueRef divisor = ctx->i32_1;
-
-			if (divisor_is_fetched) {
-				divisor = buffer_load_const(ctx, instance_divisor_constbuf,
-							    LLVMConstInt(ctx->i32, i * 4, 0));
-				divisor = ac_to_integer(&ctx->ac, divisor);
+		LLVMValueRef index = NULL;
+
+		if (divisor_is_one) {
+			index = ctx->abi.instance_id;
+		} else if (divisor_is_fetched) {
+			LLVMValueRef udiv_factors[4];
+
+			for (unsigned j = 0; j < 4; j++) {
+				udiv_factors[j] =
+					buffer_load_const(ctx, instance_divisor_constbuf,
+							  LLVMConstInt(ctx->i32, i*16 + j*4, 0));
+				udiv_factors[j] = ac_to_integer(&ctx->ac, udiv_factors[j]);
 			}
+			/* The faster NUW version doesn't work when InstanceID == UINT_MAX.
+			 * Such InstanceID might not be achievable in a reasonable time though.
+			 */
+			index = ac_build_fast_udiv_nuw(&ctx->ac, ctx->abi.instance_id,
+						       udiv_factors[0], udiv_factors[1],
+						       udiv_factors[2], udiv_factors[3]);
+		}
 
-			/* InstanceID / Divisor + StartInstance */
-			index = get_instance_index_for_fetch(ctx,
-							     user_sgpr_base +
-							     SI_SGPR_START_INSTANCE,
-							     divisor);
+		if (divisor_is_one || divisor_is_fetched) {
+			/* Add StartInstance. */
+			index = LLVMBuildAdd(ctx->ac.builder, index,
+					     LLVMGetParam(ctx->main_fn, user_sgpr_base +
+							  SI_SGPR_START_INSTANCE), "");
 		} else {
 			/* VertexID + BaseVertex */
 			index = LLVMBuildAdd(ctx->ac.builder,
@@ -7332,7 +7273,7 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx,
 }
 
 static bool si_get_vs_prolog(struct si_screen *sscreen,
-			     LLVMTargetMachineRef tm,
+			     struct ac_llvm_compiler *compiler,
 			     struct si_shader *shader,
 			     struct pipe_debug_callback *debug,
 			     struct si_shader *main_part,
@@ -7350,7 +7291,7 @@ static bool si_get_vs_prolog(struct si_screen *sscreen,
 
 	shader->prolog =
 		si_get_shader_part(sscreen, &sscreen->vs_prologs,
-				   PIPE_SHADER_VERTEX, true, &prolog_key, tm,
+				   PIPE_SHADER_VERTEX, true, &prolog_key, compiler,
 				   debug, si_build_vs_prolog_function,
 				   "Vertex Shader Prolog");
 	return shader->prolog != NULL;
@@ -7360,11 +7301,11 @@ static bool si_get_vs_prolog(struct si_screen *sscreen,
  * Select and compile (or reuse) vertex shader parts (prolog & epilog).
  */
 static bool si_shader_select_vs_parts(struct si_screen *sscreen,
-				      LLVMTargetMachineRef tm,
+				      struct ac_llvm_compiler *compiler,
 				      struct si_shader *shader,
 				      struct pipe_debug_callback *debug)
 {
-	return si_get_vs_prolog(sscreen, tm, shader, debug, shader,
+	return si_get_vs_prolog(sscreen, compiler, shader, debug, shader,
 				&shader->key.part.vs.prolog);
 }
 
@@ -7398,8 +7339,6 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx,
 		add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		add_arg(&fninfo, ARG_SGPR, ctx->i32);
-		if (!HAVE_32BIT_POINTERS)
-			add_arg(&fninfo, ARG_SGPR, ctx->ac.intptr);
 		ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		add_arg(&fninfo, ARG_SGPR, ctx->i32);
 		ctx->param_tcs_out_lds_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
@@ -7428,7 +7367,7 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx,
 
 	/* Create the function. */
 	si_create_function(ctx, "tcs_epilog", NULL, 0, &fninfo,
-			   ctx->screen->info.chip_class >= CIK ? 128 : 64);
+			   ctx->screen->info.chip_class >= GFX7 ? 128 : 64);
 	ac_declare_lds_as_pointer(&ctx->ac);
 	func = ctx->main_fn;
 
@@ -7449,7 +7388,7 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx,
  * Select and compile (or reuse) TCS parts (epilog).
  */
 static bool si_shader_select_tcs_parts(struct si_screen *sscreen,
-				       LLVMTargetMachineRef tm,
+				       struct ac_llvm_compiler *compiler,
 				       struct si_shader *shader,
 				       struct pipe_debug_callback *debug)
 {
@@ -7457,7 +7396,7 @@ static bool si_shader_select_tcs_parts(struct si_screen *sscreen,
 		struct si_shader *ls_main_part =
 			shader->key.part.tcs.ls->main_shader_part_ls;
 
-		if (!si_get_vs_prolog(sscreen, tm, shader, debug, ls_main_part,
+		if (!si_get_vs_prolog(sscreen, compiler, shader, debug, ls_main_part,
 				      &shader->key.part.tcs.ls_prolog))
 			return false;
 
@@ -7471,7 +7410,7 @@ static bool si_shader_select_tcs_parts(struct si_screen *sscreen,
 
 	shader->epilog = si_get_shader_part(sscreen, &sscreen->tcs_epilogs,
 					    PIPE_SHADER_TESS_CTRL, false,
-					    &epilog_key, tm, debug,
+					    &epilog_key, compiler, debug,
 					    si_build_tcs_epilog_function,
 					    "Tessellation Control Shader Epilog");
 	return shader->epilog != NULL;
@@ -7481,7 +7420,7 @@ static bool si_shader_select_tcs_parts(struct si_screen *sscreen,
  * Select and compile (or reuse) GS parts (prolog).
  */
 static bool si_shader_select_gs_parts(struct si_screen *sscreen,
-				      LLVMTargetMachineRef tm,
+				      struct ac_llvm_compiler *compiler,
 				      struct si_shader *shader,
 				      struct pipe_debug_callback *debug)
 {
@@ -7490,7 +7429,7 @@ static bool si_shader_select_gs_parts(struct si_screen *sscreen,
 			shader->key.part.gs.es->main_shader_part_es;
 
 		if (shader->key.part.gs.es->type == PIPE_SHADER_VERTEX &&
-		    !si_get_vs_prolog(sscreen, tm, shader, debug, es_main_part,
+		    !si_get_vs_prolog(sscreen, compiler, shader, debug, es_main_part,
 				      &shader->key.part.gs.vs_prolog))
 			return false;
 
@@ -7506,7 +7445,7 @@ static bool si_shader_select_gs_parts(struct si_screen *sscreen,
 
 	shader->prolog2 = si_get_shader_part(sscreen, &sscreen->gs_prologs,
 					    PIPE_SHADER_GEOMETRY, true,
-					    &prolog_key, tm, debug,
+					    &prolog_key, compiler, debug,
 					    si_build_gs_prolog_function,
 					    "Geometry Shader Prolog");
 	return shader->prolog2 != NULL;
@@ -7709,7 +7648,7 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx,
 							  interp_vgpr, "");
 			interp[1] = LLVMBuildExtractValue(ctx->ac.builder, ret,
 							  interp_vgpr + 1, "");
-			interp_ij = lp_build_gather_values(&ctx->gallivm, interp, 2);
+			interp_ij = ac_build_gather_values(&ctx->ac, interp, 2);
 		}
 
 		/* Use the absolute location of the input. */
@@ -7894,7 +7833,7 @@ static void si_build_ps_epilog_function(struct si_shader_context *ctx,
  * Select and compile (or reuse) pixel shader parts (prolog & epilog).
  */
 static bool si_shader_select_ps_parts(struct si_screen *sscreen,
-				      LLVMTargetMachineRef tm,
+				      struct ac_llvm_compiler *compiler,
 				      struct si_shader *shader,
 				      struct pipe_debug_callback *debug)
 {
@@ -7909,7 +7848,7 @@ static bool si_shader_select_ps_parts(struct si_screen *sscreen,
 		shader->prolog =
 			si_get_shader_part(sscreen, &sscreen->ps_prologs,
 					   PIPE_SHADER_FRAGMENT, true,
-					   &prolog_key, tm, debug,
+					   &prolog_key, compiler, debug,
 					   si_build_ps_prolog_function,
 					   "Fragment Shader Prolog");
 		if (!shader->prolog)
@@ -7922,7 +7861,7 @@ static bool si_shader_select_ps_parts(struct si_screen *sscreen,
 	shader->epilog =
 		si_get_shader_part(sscreen, &sscreen->ps_epilogs,
 				   PIPE_SHADER_FRAGMENT, false,
-				   &epilog_key, tm, debug,
+				   &epilog_key, compiler, debug,
 				   si_build_ps_epilog_function,
 				   "Fragment Shader Epilog");
 	if (!shader->epilog)
@@ -8006,8 +7945,7 @@ void si_multiwave_lds_size_workaround(struct si_screen *sscreen,
 	 *   It applies to workgroup sizes of more than one wavefront.
 	 */
 	if (sscreen->info.family == CHIP_BONAIRE ||
-	    sscreen->info.family == CHIP_KABINI ||
-	    sscreen->info.family == CHIP_MULLINS)
+	    sscreen->info.family == CHIP_KABINI)
 		*lds_size = MAX2(*lds_size, 8);
 }
 
@@ -8025,7 +7963,7 @@ static void si_fix_resource_usage(struct si_screen *sscreen,
 	}
 }
 
-int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
+bool si_shader_create(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
 		     struct si_shader *shader,
 		     struct pipe_debug_callback *debug)
 {
@@ -8043,9 +7981,9 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
 		/* Monolithic shader (compiled as a whole, has many variants,
 		 * may take a long time to compile).
 		 */
-		r = si_compile_tgsi_shader(sscreen, tm, shader, true, debug);
+		r = si_compile_tgsi_shader(sscreen, compiler, shader, debug);
 		if (r)
-			return r;
+			return false;
 	} else {
 		/* The shader consists of several parts:
 		 *
@@ -8063,7 +8001,7 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
 		 */
 
 		if (!mainp)
-			return -1;
+			return false;
 
 		/* Copy the compiled TGSI shader data over. */
 		shader->is_binary_shared = true;
@@ -8083,22 +8021,22 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
 		/* Select prologs and/or epilogs. */
 		switch (sel->type) {
 		case PIPE_SHADER_VERTEX:
-			if (!si_shader_select_vs_parts(sscreen, tm, shader, debug))
-				return -1;
+			if (!si_shader_select_vs_parts(sscreen, compiler, shader, debug))
+				return false;
 			break;
 		case PIPE_SHADER_TESS_CTRL:
-			if (!si_shader_select_tcs_parts(sscreen, tm, shader, debug))
-				return -1;
+			if (!si_shader_select_tcs_parts(sscreen, compiler, shader, debug))
+				return false;
 			break;
 		case PIPE_SHADER_TESS_EVAL:
 			break;
 		case PIPE_SHADER_GEOMETRY:
-			if (!si_shader_select_gs_parts(sscreen, tm, shader, debug))
-				return -1;
+			if (!si_shader_select_gs_parts(sscreen, compiler, shader, debug))
+				return false;
 			break;
 		case PIPE_SHADER_FRAGMENT:
-			if (!si_shader_select_ps_parts(sscreen, tm, shader, debug))
-				return -1;
+			if (!si_shader_select_ps_parts(sscreen, compiler, shader, debug))
+				return false;
 
 			/* Make sure we have at least as many VGPRs as there
 			 * are allocated inputs.
@@ -8126,9 +8064,9 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
 			shader->config.spilled_vgprs =
 				MAX2(shader->config.spilled_vgprs,
 				     shader->previous_stage->config.spilled_vgprs);
-			shader->config.private_mem_vgprs =
-				MAX2(shader->config.private_mem_vgprs,
-				     shader->previous_stage->config.private_mem_vgprs);
+			shader->info.private_mem_vgprs =
+				MAX2(shader->info.private_mem_vgprs,
+				     shader->previous_stage->info.private_mem_vgprs);
 			shader->config.scratch_bytes_per_wave =
 				MAX2(shader->config.scratch_bytes_per_wave,
 				     shader->previous_stage->config.scratch_bytes_per_wave);
@@ -8150,29 +8088,31 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm,
 		si_calculate_max_simd_waves(shader);
 	}
 
+	if (sscreen->info.chip_class >= GFX9 && sel->type == PIPE_SHADER_GEOMETRY)
+		gfx9_get_gs_info(shader->previous_stage_sel, sel, &shader->gs_info);
+
 	si_fix_resource_usage(sscreen, shader);
 	si_shader_dump(sscreen, shader, debug, sel->info.processor,
 		       stderr, true);
 
 	/* Upload. */
-	r = si_shader_binary_upload(sscreen, shader);
-	if (r) {
+	if (!si_shader_binary_upload(sscreen, shader, 0)) {
 		fprintf(stderr, "LLVM failed to upload shader\n");
-		return r;
+		return false;
 	}
 
-	return 0;
+	return true;
 }
 
 void si_shader_destroy(struct si_shader *shader)
 {
 	if (shader->scratch_bo)
-		r600_resource_reference(&shader->scratch_bo, NULL);
+		si_resource_reference(&shader->scratch_bo, NULL);
 
-	r600_resource_reference(&shader->bo, NULL);
+	si_resource_reference(&shader->bo, NULL);
 
 	if (!shader->is_binary_shared)
-		ac_shader_binary_clean(&shader->binary);
+		si_shader_binary_clean(&shader->binary);
 
 	free(shader->shader_log);
 }