From: Marek Olšák <marek.olsak@amd.com>
Date: Mon, 30 Dec 2019 19:23:16 +0000 (-0500)
Subject: radeonsi/gfx10: implement NGG culling for 4x wave32 subgroups
X-Git-Url: https://git.libre-soc.org/?p=mesa.git;a=commitdiff_plain;h=8db00a51f85109e958631ef74a458b0614f37097

radeonsi/gfx10: implement NGG culling for 4x wave32 subgroups

Acked-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
---

diff --git a/src/gallium/drivers/radeonsi/gfx10_shader_ngg.c b/src/gallium/drivers/radeonsi/gfx10_shader_ngg.c
index a25c89bac56..8092b796b5d 100644
--- a/src/gallium/drivers/radeonsi/gfx10_shader_ngg.c
+++ b/src/gallium/drivers/radeonsi/gfx10_shader_ngg.c
@@ -28,6 +28,7 @@
 
 #include "util/u_memory.h"
 #include "util/u_prim.h"
+#include "ac_llvm_cull.h"
 
 static LLVMValueRef get_wave_id_in_tg(struct si_shader_context *ctx)
 {
@@ -141,14 +142,44 @@ void gfx10_ngg_build_sendmsg_gs_alloc_req(struct si_shader_context *ctx)
 }
 
 void gfx10_ngg_build_export_prim(struct si_shader_context *ctx,
-				 LLVMValueRef user_edgeflags[3])
+				 LLVMValueRef user_edgeflags[3],
+				 LLVMValueRef prim_passthrough)
 {
-	if (gfx10_is_ngg_passthrough(ctx->shader)) {
+	LLVMBuilderRef builder = ctx->ac.builder;
+
+	if (gfx10_is_ngg_passthrough(ctx->shader) ||
+	    ctx->shader->key.opt.ngg_culling) {
 		ac_build_ifcc(&ctx->ac, si_is_gs_thread(ctx), 6001);
 		{
 			struct ac_ngg_prim prim = {};
 
-			prim.passthrough = ac_get_arg(&ctx->ac, ctx->gs_vtx01_offset);
+			if (prim_passthrough)
+				prim.passthrough = prim_passthrough;
+			else
+				prim.passthrough = ac_get_arg(&ctx->ac, ctx->gs_vtx01_offset);
+
+			/* This is only used with NGG culling, which returns the NGG
+			 * passthrough prim export encoding.
+			 */
+			if (ctx->shader->selector->info.writes_edgeflag) {
+				unsigned all_bits_no_edgeflags = ~SI_NGG_PRIM_EDGE_FLAG_BITS;
+				LLVMValueRef edgeflags = LLVMConstInt(ctx->i32, all_bits_no_edgeflags, 0);
+
+				unsigned num_vertices;
+				ngg_get_vertices_per_prim(ctx, &num_vertices);
+
+				for (unsigned i = 0; i < num_vertices; i++) {
+					unsigned shift = 9 + i*10;
+					LLVMValueRef edge;
+
+					edge = LLVMBuildLoad(builder, user_edgeflags[i], "");
+					edge = LLVMBuildZExt(builder, edge, ctx->i32, "");
+					edge = LLVMBuildShl(builder, edge, LLVMConstInt(ctx->i32, shift, 0), "");
+					edgeflags = LLVMBuildOr(builder, edgeflags, edge, "");
+				}
+				prim.passthrough = LLVMBuildAnd(builder, prim.passthrough, edgeflags, "");
+			}
+
 			ac_build_export_prim(&ctx->ac, &prim);
 		}
 		ac_build_endif(&ctx->ac, 6001);
@@ -535,6 +566,51 @@ static void build_streamout(struct si_shader_context *ctx,
 	}
 }
 
+/* LDS layout of ES vertex data for NGG culling. */
+enum {
+	/* Byte 0: Boolean ES thread accepted (unculled) flag, and later the old
+	 *         ES thread ID. After vertex compaction, compacted ES threads
+	 *         store the old thread ID here to copy input VGPRs from uncompacted
+	 *         ES threads.
+	 * Byte 1: New ES thread ID, loaded by GS to prepare the prim export value.
+	 * Byte 2: TES rel patch ID
+	 * Byte 3: Unused
+	 */
+	lds_byte0_accept_flag = 0,
+	lds_byte0_old_thread_id = 0,
+	lds_byte1_new_thread_id,
+	lds_byte2_tes_rel_patch_id,
+	lds_byte3_unused,
+
+	lds_packed_data = 0, /* lds_byteN_... */
+
+	lds_pos_x,
+	lds_pos_y,
+	lds_pos_z,
+	lds_pos_w,
+	lds_pos_x_div_w,
+	lds_pos_y_div_w,
+	/* If VS: */
+	lds_vertex_id,
+	lds_instance_id, /* optional */
+	/* If TES: */
+	lds_tes_u = lds_vertex_id,
+	lds_tes_v = lds_instance_id,
+	lds_tes_patch_id, /* optional */
+};
+
+static LLVMValueRef si_build_gep_i8(struct si_shader_context *ctx,
+				    LLVMValueRef ptr, unsigned byte_index)
+{
+	assert(byte_index < 4);
+	LLVMTypeRef pi8 = LLVMPointerType(ctx->i8, AC_ADDR_SPACE_LDS);
+	LLVMValueRef index = LLVMConstInt(ctx->i32, byte_index, 0);
+
+	return LLVMBuildGEP(ctx->ac.builder,
+			    LLVMBuildPointerCast(ctx->ac.builder, ptr, pi8, ""),
+			    &index, 1, "");
+}
+
 static unsigned ngg_nogs_vertex_size(struct si_shader *shader)
 {
 	unsigned lds_vertex_size = 0;
@@ -555,6 +631,24 @@ static unsigned ngg_nogs_vertex_size(struct si_shader *shader)
 	    shader->key.mono.u.vs_export_prim_id)
 		lds_vertex_size = MAX2(lds_vertex_size, 1);
 
+	if (shader->key.opt.ngg_culling) {
+		if (shader->selector->type == PIPE_SHADER_VERTEX) {
+			STATIC_ASSERT(lds_instance_id + 1 == 9);
+			lds_vertex_size = MAX2(lds_vertex_size, 9);
+		} else {
+			assert(shader->selector->type == PIPE_SHADER_TESS_EVAL);
+
+			if (shader->selector->info.uses_primid ||
+			    shader->key.mono.u.vs_export_prim_id) {
+				STATIC_ASSERT(lds_tes_patch_id + 2 == 11);
+				lds_vertex_size = MAX2(lds_vertex_size, 11);
+			} else {
+				STATIC_ASSERT(lds_tes_v + 1 == 9);
+				lds_vertex_size = MAX2(lds_vertex_size, 9);
+			}
+		}
+	}
+
 	return lds_vertex_size;
 }
 
@@ -573,6 +667,540 @@ static LLVMValueRef ngg_nogs_vertex_ptr(struct si_shader_context *ctx,
 	return LLVMBuildGEP(ctx->ac.builder, tmp, &vtxid, 1, "");
 }
 
+static void load_bitmasks_2x64(struct si_shader_context *ctx,
+			       LLVMValueRef lds_ptr, unsigned dw_offset,
+			       LLVMValueRef mask[2], LLVMValueRef *total_bitcount)
+{
+	LLVMBuilderRef builder = ctx->ac.builder;
+	LLVMValueRef ptr64 = LLVMBuildPointerCast(builder, lds_ptr,
+						  LLVMPointerType(LLVMArrayType(ctx->i64, 2),
+								  AC_ADDR_SPACE_LDS), "");
+	for (unsigned i = 0; i < 2; i++) {
+		LLVMValueRef index = LLVMConstInt(ctx->i32, dw_offset / 2 + i, 0);
+		mask[i] = LLVMBuildLoad(builder, ac_build_gep0(&ctx->ac, ptr64, index), "");
+	}
+
+	/* We get better code if we don't use the 128-bit bitcount. */
+	*total_bitcount = LLVMBuildAdd(builder, ac_build_bit_count(&ctx->ac, mask[0]),
+				       ac_build_bit_count(&ctx->ac, mask[1]), "");
+}
+
+/**
+ * Given a total thread count, update total and per-wave thread counts in input SGPRs
+ * and return the per-wave thread count.
+ *
+ * \param new_num_threads    Total thread count on the input, per-wave thread count on the output.
+ * \param tg_info	     tg_info SGPR value
+ * \param tg_info_num_bits   the bit size of thread count field in tg_info
+ * \param tg_info_shift      the bit offset of the thread count field in tg_info
+ * \param wave_info          merged_wave_info SGPR value
+ * \param wave_info_num_bits the bit size of thread count field in merged_wave_info
+ * \param wave_info_shift    the bit offset of the thread count field in merged_wave_info
+ */
+static void update_thread_counts(struct si_shader_context *ctx,
+				 LLVMValueRef *new_num_threads,
+				 LLVMValueRef *tg_info,
+				 unsigned tg_info_num_bits,
+				 unsigned tg_info_shift,
+				 LLVMValueRef *wave_info,
+				 unsigned wave_info_num_bits,
+				 unsigned wave_info_shift)
+{
+	LLVMBuilderRef builder = ctx->ac.builder;
+
+	/* Update the total thread count. */
+	unsigned tg_info_mask = ~(u_bit_consecutive(0, tg_info_num_bits) << tg_info_shift);
+	*tg_info = LLVMBuildAnd(builder, *tg_info,
+				LLVMConstInt(ctx->i32, tg_info_mask, 0), "");
+	*tg_info = LLVMBuildOr(builder, *tg_info,
+			       LLVMBuildShl(builder, *new_num_threads,
+					    LLVMConstInt(ctx->i32, tg_info_shift, 0), ""), "");
+
+	/* Update the per-wave thread count. */
+	LLVMValueRef prev_threads = LLVMBuildMul(builder, get_wave_id_in_tg(ctx),
+						 LLVMConstInt(ctx->i32, ctx->ac.wave_size, 0), "");
+	*new_num_threads = LLVMBuildSub(builder, *new_num_threads, prev_threads, "");
+	*new_num_threads = ac_build_imax(&ctx->ac, *new_num_threads, ctx->i32_0);
+	*new_num_threads = ac_build_imin(&ctx->ac, *new_num_threads,
+					LLVMConstInt(ctx->i32, ctx->ac.wave_size, 0));
+	unsigned wave_info_mask = ~(u_bit_consecutive(0, wave_info_num_bits) << wave_info_shift);
+	*wave_info = LLVMBuildAnd(builder, *wave_info,
+				  LLVMConstInt(ctx->i32, wave_info_mask, 0), "");
+	*wave_info = LLVMBuildOr(builder, *wave_info,
+				 LLVMBuildShl(builder, *new_num_threads,
+					      LLVMConstInt(ctx->i32, wave_info_shift, 0), ""), "");
+}
+
+/**
+ * Cull primitives for NGG VS or TES, then compact vertices, which happens
+ * before the VS or TES main function. Return values for the main function.
+ * Also return the position, which is passed to the shader as an input,
+ * so that we don't compute it twice.
+ */
+void gfx10_emit_ngg_culling_epilogue_4x_wave32(struct ac_shader_abi *abi,
+					       unsigned max_outputs,
+					       LLVMValueRef *addrs)
+{
+	struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+	struct si_shader *shader = ctx->shader;
+	struct si_shader_selector *sel = shader->selector;
+	struct si_shader_info *info = &sel->info;
+	LLVMBuilderRef builder = ctx->ac.builder;
+
+	assert(shader->key.opt.ngg_culling);
+	assert(shader->key.as_ngg);
+	assert(sel->type == PIPE_SHADER_VERTEX ||
+	       (sel->type == PIPE_SHADER_TESS_EVAL && !shader->key.as_es));
+
+	LLVMValueRef position[4] = {};
+	for (unsigned i = 0; i < info->num_outputs; i++) {
+		switch (info->output_semantic_name[i]) {
+		case TGSI_SEMANTIC_POSITION:
+			for (unsigned j = 0; j < 4; j++) {
+				position[j] = LLVMBuildLoad(ctx->ac.builder,
+							    addrs[4 * i + j], "");
+			}
+			break;
+		}
+	}
+	assert(position[0]);
+
+	/* Store Position.XYZW into LDS. */
+	LLVMValueRef es_vtxptr = ngg_nogs_vertex_ptr(ctx, get_thread_id_in_tg(ctx));
+	for (unsigned chan = 0; chan < 4; chan++) {
+		LLVMBuildStore(builder, ac_to_integer(&ctx->ac, position[chan]),
+				ac_build_gep0(&ctx->ac, es_vtxptr,
+					      LLVMConstInt(ctx->i32, lds_pos_x + chan, 0)));
+	}
+	/* Store Position.XY / W into LDS. */
+	for (unsigned chan = 0; chan < 2; chan++) {
+		LLVMValueRef val = ac_build_fdiv(&ctx->ac, position[chan], position[3]);
+		LLVMBuildStore(builder, ac_to_integer(&ctx->ac, val),
+				ac_build_gep0(&ctx->ac, es_vtxptr,
+					      LLVMConstInt(ctx->i32, lds_pos_x_div_w + chan, 0)));
+	}
+
+	/* Store VertexID and InstanceID. ES threads will have to load them
+	 * from LDS after vertex compaction and use them instead of their own
+	 * system values.
+	 */
+	bool uses_instance_id = false;
+	bool uses_tes_prim_id = false;
+	LLVMValueRef packed_data = ctx->i32_0;
+
+	if (ctx->type == PIPE_SHADER_VERTEX) {
+		uses_instance_id = sel->info.uses_instanceid ||
+				   shader->key.part.vs.prolog.instance_divisor_is_one ||
+				   shader->key.part.vs.prolog.instance_divisor_is_fetched;
+
+		LLVMBuildStore(builder, ctx->abi.vertex_id,
+			       ac_build_gep0(&ctx->ac, es_vtxptr,
+					     LLVMConstInt(ctx->i32, lds_vertex_id, 0)));
+		if (uses_instance_id) {
+			LLVMBuildStore(builder, ctx->abi.instance_id,
+				       ac_build_gep0(&ctx->ac, es_vtxptr,
+						     LLVMConstInt(ctx->i32, lds_instance_id, 0)));
+		}
+	} else {
+		uses_tes_prim_id = sel->info.uses_primid ||
+				   shader->key.mono.u.vs_export_prim_id;
+
+		assert(ctx->type == PIPE_SHADER_TESS_EVAL);
+		LLVMBuildStore(builder, ac_to_integer(&ctx->ac, ac_get_arg(&ctx->ac, ctx->tes_u)),
+			       ac_build_gep0(&ctx->ac, es_vtxptr,
+					     LLVMConstInt(ctx->i32, lds_tes_u, 0)));
+		LLVMBuildStore(builder, ac_to_integer(&ctx->ac, ac_get_arg(&ctx->ac, ctx->tes_v)),
+			       ac_build_gep0(&ctx->ac, es_vtxptr,
+					     LLVMConstInt(ctx->i32, lds_tes_v, 0)));
+		packed_data = LLVMBuildShl(builder, ac_get_arg(&ctx->ac, ctx->tes_rel_patch_id),
+					   LLVMConstInt(ctx->i32, lds_byte2_tes_rel_patch_id * 8, 0), "");
+		if (uses_tes_prim_id) {
+			LLVMBuildStore(builder, ac_get_arg(&ctx->ac, ctx->args.tes_patch_id),
+				       ac_build_gep0(&ctx->ac, es_vtxptr,
+						     LLVMConstInt(ctx->i32, lds_tes_patch_id, 0)));
+		}
+	}
+	/* Initialize the packed data. */
+	LLVMBuildStore(builder, packed_data,
+		       ac_build_gep0(&ctx->ac, es_vtxptr,
+				     LLVMConstInt(ctx->i32, lds_packed_data, 0)));
+	ac_build_endif(&ctx->ac, ctx->merged_wrap_if_label);
+
+	LLVMValueRef tid = ac_get_thread_id(&ctx->ac);
+
+	/* Initialize the last 3 gs_ngg_scratch dwords to 0, because we may have less
+	 * than 4 waves, but we always read all 4 values. This is where the thread
+	 * bitmasks of unculled threads will be stored.
+	 *
+	 * gs_ngg_scratch layout: esmask[0..3]
+	 */
+	ac_build_ifcc(&ctx->ac,
+		      LLVMBuildICmp(builder, LLVMIntULT, get_thread_id_in_tg(ctx),
+				    LLVMConstInt(ctx->i32, 3, 0), ""), 16101);
+	{
+		LLVMValueRef index = LLVMBuildAdd(builder, tid, ctx->i32_1, "");
+		LLVMBuildStore(builder, ctx->i32_0,
+			       ac_build_gep0(&ctx->ac, ctx->gs_ngg_scratch, index));
+	}
+	ac_build_endif(&ctx->ac, 16101);
+	ac_build_s_barrier(&ctx->ac);
+
+	/* The hardware requires that there are no holes between unculled vertices,
+	 * which means we have to pack ES threads, i.e. reduce the ES thread count
+	 * and move ES input VGPRs to lower threads. The upside is that varyings
+	 * are only fetched and computed for unculled vertices.
+	 *
+	 * Vertex compaction in GS threads:
+	 *
+	 * Part 1: Compute the surviving vertex mask in GS threads:
+	 * - Compute 4 32-bit surviving vertex masks in LDS. (max 4 waves)
+	 *   - In GS, notify ES threads whether the vertex survived.
+	 *   - Barrier
+	 *   - ES threads will create the mask and store it in LDS.
+	 * - Barrier
+	 * - Each GS thread loads the vertex masks from LDS.
+	 *
+	 * Part 2: Compact ES threads in GS threads:
+	 * - Compute the prefix sum for all 3 vertices from the masks. These are the new
+	 *   thread IDs for each vertex within the primitive.
+	 * - Write the value of the old thread ID into the LDS address of the new thread ID.
+	 *   The ES thread will load the old thread ID and use it to load the position, VertexID,
+	 *   and InstanceID.
+	 * - Update vertex indices and null flag in the GS input VGPRs.
+	 * - Barrier
+	 *
+	 * Part 3: Update inputs GPRs
+	 * - For all waves, update per-wave thread counts in input SGPRs.
+	 * - In ES threads, update the ES input VGPRs (VertexID, InstanceID, TES inputs).
+	 */
+
+	LLVMValueRef vtxindex[] = {
+		si_unpack_param(ctx, ctx->gs_vtx01_offset, 0, 16),
+		si_unpack_param(ctx, ctx->gs_vtx01_offset, 16, 16),
+		si_unpack_param(ctx, ctx->gs_vtx23_offset, 0, 16),
+	};
+	LLVMValueRef gs_vtxptr[] = {
+		ngg_nogs_vertex_ptr(ctx, vtxindex[0]),
+		ngg_nogs_vertex_ptr(ctx, vtxindex[1]),
+		ngg_nogs_vertex_ptr(ctx, vtxindex[2]),
+	};
+	es_vtxptr = ngg_nogs_vertex_ptr(ctx, get_thread_id_in_tg(ctx));
+
+	LLVMValueRef gs_accepted = ac_build_alloca(&ctx->ac, ctx->i32, "");
+
+	/* Do culling in GS threads. */
+	ac_build_ifcc(&ctx->ac, si_is_gs_thread(ctx), 16002);
+	{
+		/* Load positions. */
+		LLVMValueRef pos[3][4] = {};
+		for (unsigned vtx = 0; vtx < 3; vtx++) {
+			for (unsigned chan = 0; chan < 4; chan++) {
+				unsigned index;
+				if (chan == 0 || chan == 1)
+					index = lds_pos_x_div_w + chan;
+				else if (chan == 3)
+					index = lds_pos_w;
+				else
+					continue;
+
+				LLVMValueRef addr = ac_build_gep0(&ctx->ac, gs_vtxptr[vtx],
+								  LLVMConstInt(ctx->i32, index, 0));
+				pos[vtx][chan] = LLVMBuildLoad(builder, addr, "");
+				pos[vtx][chan] = ac_to_float(&ctx->ac, pos[vtx][chan]);
+			}
+		}
+
+		/* Load the viewport state for small prim culling. */
+		LLVMValueRef vp = ac_build_load_invariant(&ctx->ac,
+							  ac_get_arg(&ctx->ac, ctx->small_prim_cull_info),
+							  ctx->i32_0);
+		vp = LLVMBuildBitCast(builder, vp, ctx->v4f32, "");
+		LLVMValueRef vp_scale[2], vp_translate[2];
+		vp_scale[0] = ac_llvm_extract_elem(&ctx->ac, vp, 0);
+		vp_scale[1] = ac_llvm_extract_elem(&ctx->ac, vp, 1);
+		vp_translate[0] = ac_llvm_extract_elem(&ctx->ac, vp, 2);
+		vp_translate[1] = ac_llvm_extract_elem(&ctx->ac, vp, 3);
+
+		/* Get the small prim filter precision. */
+		LLVMValueRef small_prim_precision = si_unpack_param(ctx, ctx->vs_state_bits, 7, 4);
+		small_prim_precision = LLVMBuildOr(builder, small_prim_precision,
+						   LLVMConstInt(ctx->i32, 0x70, 0), "");
+		small_prim_precision = LLVMBuildShl(builder, small_prim_precision,
+						    LLVMConstInt(ctx->i32, 23, 0), "");
+		small_prim_precision = LLVMBuildBitCast(builder, small_prim_precision, ctx->f32, "");
+
+		/* Execute culling code. */
+		struct ac_cull_options options = {};
+		options.cull_front = shader->key.opt.ngg_culling & SI_NGG_CULL_FRONT_FACE;
+		options.cull_back = shader->key.opt.ngg_culling & SI_NGG_CULL_BACK_FACE;
+		options.cull_view_xy = shader->key.opt.ngg_culling & SI_NGG_CULL_VIEW_SMALLPRIMS;
+		options.cull_small_prims = options.cull_view_xy;
+		options.cull_zero_area = options.cull_front || options.cull_back;
+		options.cull_w = true;
+
+		/* Tell ES threads whether their vertex survived. */
+		ac_build_ifcc(&ctx->ac, ac_cull_triangle(&ctx->ac, pos, ctx->i1true,
+							 vp_scale, vp_translate,
+							 small_prim_precision, &options), 16003);
+		{
+			LLVMBuildStore(builder, ctx->ac.i32_1, gs_accepted);
+			for (unsigned vtx = 0; vtx < 3; vtx++) {
+				LLVMBuildStore(builder, ctx->ac.i8_1,
+					       si_build_gep_i8(ctx, gs_vtxptr[vtx], lds_byte0_accept_flag));
+			}
+		}
+		ac_build_endif(&ctx->ac, 16003);
+	}
+	ac_build_endif(&ctx->ac, 16002);
+	ac_build_s_barrier(&ctx->ac);
+
+	gs_accepted = LLVMBuildLoad(builder, gs_accepted, "");
+
+	LLVMValueRef es_accepted = ac_build_alloca(&ctx->ac, ctx->i1, "");
+
+	/* Convert the per-vertex flag to a thread bitmask in ES threads and store it in LDS. */
+	ac_build_ifcc(&ctx->ac, si_is_es_thread(ctx), 16007);
+	{
+		LLVMValueRef es_accepted_flag =
+			LLVMBuildLoad(builder,
+				      si_build_gep_i8(ctx, es_vtxptr, lds_byte0_accept_flag), "");
+
+		LLVMValueRef es_accepted_bool = LLVMBuildICmp(builder, LLVMIntNE,
+							      es_accepted_flag, ctx->ac.i8_0, "");
+		LLVMValueRef es_mask = ac_get_i1_sgpr_mask(&ctx->ac, es_accepted_bool);
+
+		LLVMBuildStore(builder, es_accepted_bool, es_accepted);
+
+		ac_build_ifcc(&ctx->ac, LLVMBuildICmp(builder, LLVMIntEQ,
+						      tid, ctx->i32_0, ""), 16008);
+		{
+			LLVMBuildStore(builder, es_mask,
+				       ac_build_gep0(&ctx->ac, ctx->gs_ngg_scratch,
+						     get_wave_id_in_tg(ctx)));
+		}
+		ac_build_endif(&ctx->ac, 16008);
+	}
+	ac_build_endif(&ctx->ac, 16007);
+	ac_build_s_barrier(&ctx->ac);
+
+	/* Load the vertex masks and compute the new ES thread count. */
+	LLVMValueRef es_mask[2], new_num_es_threads, kill_wave;
+	load_bitmasks_2x64(ctx, ctx->gs_ngg_scratch, 0, es_mask, &new_num_es_threads);
+	new_num_es_threads = ac_build_readlane_no_opt_barrier(&ctx->ac, new_num_es_threads, NULL);
+
+	/* ES threads compute their prefix sum, which is the new ES thread ID.
+	 * Then they write the value of the old thread ID into the LDS address
+	 * of the new thread ID. It will be used it to load input VGPRs from
+	 * the old thread's LDS location.
+	 */
+	ac_build_ifcc(&ctx->ac, LLVMBuildLoad(builder, es_accepted, ""), 16009);
+	{
+		LLVMValueRef old_id = get_thread_id_in_tg(ctx);
+		LLVMValueRef new_id = ac_prefix_bitcount_2x64(&ctx->ac, es_mask, old_id);
+
+		LLVMBuildStore(builder, LLVMBuildTrunc(builder, old_id, ctx->i8, ""),
+			       si_build_gep_i8(ctx, ngg_nogs_vertex_ptr(ctx, new_id),
+					       lds_byte0_old_thread_id));
+		LLVMBuildStore(builder, LLVMBuildTrunc(builder, new_id, ctx->i8, ""),
+			       si_build_gep_i8(ctx, es_vtxptr, lds_byte1_new_thread_id));
+	}
+	ac_build_endif(&ctx->ac, 16009);
+
+	/* Kill waves that have inactive threads. */
+	kill_wave = LLVMBuildICmp(builder, LLVMIntULE,
+				  ac_build_imax(&ctx->ac, new_num_es_threads, ngg_get_prim_cnt(ctx)),
+				  LLVMBuildMul(builder, get_wave_id_in_tg(ctx),
+					       LLVMConstInt(ctx->i32, ctx->ac.wave_size, 0), ""), "");
+	ac_build_ifcc(&ctx->ac, kill_wave, 19202);
+	{
+		/* If we are killing wave 0, send that there are no primitives
+		 * in this threadgroup.
+		 */
+		ac_build_sendmsg_gs_alloc_req(&ctx->ac, get_wave_id_in_tg(ctx),
+					      ctx->i32_0, ctx->i32_0);
+		ac_build_s_endpgm(&ctx->ac);
+	}
+	ac_build_endif(&ctx->ac, 19202);
+	ac_build_s_barrier(&ctx->ac);
+
+	/* Send the final vertex and primitive counts. */
+	ac_build_sendmsg_gs_alloc_req(&ctx->ac, get_wave_id_in_tg(ctx),
+				      new_num_es_threads, ngg_get_prim_cnt(ctx));
+
+	/* Update thread counts in SGPRs. */
+	LLVMValueRef new_gs_tg_info = ac_get_arg(&ctx->ac, ctx->gs_tg_info);
+	LLVMValueRef new_merged_wave_info = ac_get_arg(&ctx->ac, ctx->merged_wave_info);
+
+	/* This also converts the thread count from the total count to the per-wave count. */
+	update_thread_counts(ctx, &new_num_es_threads, &new_gs_tg_info, 9, 12,
+			     &new_merged_wave_info, 8, 0);
+
+	/* Update vertex indices in VGPR0 (same format as NGG passthrough). */
+	LLVMValueRef new_vgpr0 = ac_build_alloca_undef(&ctx->ac, ctx->i32, "");
+
+	/* Set the null flag at the beginning (culled), and then
+	 * overwrite it for accepted primitives.
+	 */
+	LLVMBuildStore(builder, LLVMConstInt(ctx->i32, 1u << 31, 0), new_vgpr0);
+
+	/* Get vertex indices after vertex compaction. */
+	ac_build_ifcc(&ctx->ac, LLVMBuildTrunc(builder, gs_accepted, ctx->i1, ""), 16011);
+	{
+		struct ac_ngg_prim prim = {};
+		prim.num_vertices = 3;
+		prim.isnull = ctx->i1false;
+
+		for (unsigned vtx = 0; vtx < 3; vtx++) {
+			prim.index[vtx] =
+				LLVMBuildLoad(builder,
+					      si_build_gep_i8(ctx, gs_vtxptr[vtx],
+							      lds_byte1_new_thread_id), "");
+			prim.index[vtx] = LLVMBuildZExt(builder, prim.index[vtx], ctx->i32, "");
+			prim.edgeflag[vtx] = ngg_get_initial_edgeflag(ctx, vtx);
+		}
+
+		/* Set the new GS input VGPR. */
+		LLVMBuildStore(builder, ac_pack_prim_export(&ctx->ac, &prim), new_vgpr0);
+	}
+	ac_build_endif(&ctx->ac, 16011);
+
+	if (gfx10_ngg_export_prim_early(shader))
+		gfx10_ngg_build_export_prim(ctx, NULL, LLVMBuildLoad(builder, new_vgpr0, ""));
+
+	/* Set the new ES input VGPRs. */
+	LLVMValueRef es_data[4];
+	LLVMValueRef old_thread_id = ac_build_alloca_undef(&ctx->ac, ctx->i32, "");
+
+	for (unsigned i = 0; i < 4; i++)
+		es_data[i] = ac_build_alloca_undef(&ctx->ac, ctx->i32, "");
+
+	ac_build_ifcc(&ctx->ac, LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, tid,
+					      new_num_es_threads, ""), 16012);
+	{
+		LLVMValueRef old_id, old_es_vtxptr, tmp;
+
+		/* Load ES input VGPRs from the ES thread before compaction. */
+		old_id = LLVMBuildLoad(builder,
+				       si_build_gep_i8(ctx, es_vtxptr, lds_byte0_old_thread_id), "");
+		old_id = LLVMBuildZExt(builder, old_id, ctx->i32, "");
+
+		LLVMBuildStore(builder, old_id, old_thread_id);
+		old_es_vtxptr = ngg_nogs_vertex_ptr(ctx, old_id);
+
+		for (unsigned i = 0; i < 2; i++) {
+			tmp = LLVMBuildLoad(builder,
+					    ac_build_gep0(&ctx->ac, old_es_vtxptr,
+							  LLVMConstInt(ctx->i32, lds_vertex_id + i, 0)), "");
+			LLVMBuildStore(builder, tmp, es_data[i]);
+		}
+
+		if (ctx->type == PIPE_SHADER_TESS_EVAL) {
+			tmp = LLVMBuildLoad(builder,
+					    si_build_gep_i8(ctx, old_es_vtxptr,
+							    lds_byte2_tes_rel_patch_id), "");
+			tmp = LLVMBuildZExt(builder, tmp, ctx->i32, "");
+			LLVMBuildStore(builder, tmp, es_data[2]);
+
+			if (uses_tes_prim_id) {
+				tmp = LLVMBuildLoad(builder,
+						    ac_build_gep0(&ctx->ac, old_es_vtxptr,
+								  LLVMConstInt(ctx->i32, lds_tes_patch_id, 0)), "");
+				LLVMBuildStore(builder, tmp, es_data[3]);
+			}
+		}
+	}
+	ac_build_endif(&ctx->ac, 16012);
+
+	/* Return values for the main function. */
+	LLVMValueRef ret = ctx->return_value;
+	LLVMValueRef val;
+
+	ret = LLVMBuildInsertValue(ctx->ac.builder, ret, new_gs_tg_info, 2, "");
+	ret = LLVMBuildInsertValue(ctx->ac.builder, ret, new_merged_wave_info, 3, "");
+	if (ctx->type == PIPE_SHADER_TESS_EVAL)
+		ret = si_insert_input_ret(ctx, ret, ctx->tcs_offchip_offset, 4);
+
+	ret = si_insert_input_ptr(ctx, ret, ctx->rw_buffers,
+				  8 + SI_SGPR_RW_BUFFERS);
+	ret = si_insert_input_ptr(ctx, ret,
+				  ctx->bindless_samplers_and_images,
+				  8 + SI_SGPR_BINDLESS_SAMPLERS_AND_IMAGES);
+	ret = si_insert_input_ptr(ctx, ret,
+				  ctx->const_and_shader_buffers,
+				  8 + SI_SGPR_CONST_AND_SHADER_BUFFERS);
+	ret = si_insert_input_ptr(ctx, ret,
+				  ctx->samplers_and_images,
+				  8 + SI_SGPR_SAMPLERS_AND_IMAGES);
+	ret = si_insert_input_ptr(ctx, ret, ctx->vs_state_bits,
+				  8 + SI_SGPR_VS_STATE_BITS);
+
+	if (ctx->type == PIPE_SHADER_VERTEX) {
+		ret = si_insert_input_ptr(ctx, ret, ctx->args.base_vertex,
+					  8 + SI_SGPR_BASE_VERTEX);
+		ret = si_insert_input_ptr(ctx, ret, ctx->args.start_instance,
+					  8 + SI_SGPR_START_INSTANCE);
+		ret = si_insert_input_ptr(ctx, ret, ctx->args.draw_id,
+					  8 + SI_SGPR_DRAWID);
+		ret = si_insert_input_ptr(ctx, ret, ctx->vertex_buffers,
+					  8 + SI_VS_NUM_USER_SGPR);
+	} else {
+		assert(ctx->type == PIPE_SHADER_TESS_EVAL);
+		ret = si_insert_input_ptr(ctx, ret, ctx->tcs_offchip_layout,
+					  8 + SI_SGPR_TES_OFFCHIP_LAYOUT);
+		ret = si_insert_input_ptr(ctx, ret, ctx->tes_offchip_addr,
+					  8 + SI_SGPR_TES_OFFCHIP_ADDR);
+	}
+
+	unsigned vgpr;
+	if (ctx->type == PIPE_SHADER_VERTEX)
+		vgpr = 8 + GFX9_VSGS_NUM_USER_SGPR + 1;
+	else
+		vgpr = 8 + GFX9_TESGS_NUM_USER_SGPR;
+
+	val = LLVMBuildLoad(builder, new_vgpr0, "");
+	ret = LLVMBuildInsertValue(builder, ret, ac_to_float(&ctx->ac, val),
+				   vgpr++, "");
+	vgpr++; /* gs_vtx23_offset */
+
+	ret = si_insert_input_ret_float(ctx, ret, ctx->args.gs_prim_id, vgpr++);
+	ret = si_insert_input_ret_float(ctx, ret, ctx->args.gs_invocation_id, vgpr++);
+	vgpr++; /* gs_vtx45_offset */
+
+	if (ctx->type == PIPE_SHADER_VERTEX) {
+		val = LLVMBuildLoad(builder, es_data[0], "");
+		ret = LLVMBuildInsertValue(builder, ret, ac_to_float(&ctx->ac, val),
+					   vgpr++, ""); /* VGPR5 - VertexID */
+		vgpr += 2;
+		if (uses_instance_id) {
+			val = LLVMBuildLoad(builder, es_data[1], "");
+			ret = LLVMBuildInsertValue(builder, ret, ac_to_float(&ctx->ac, val),
+						   vgpr++, ""); /* VGPR8 - InstanceID */
+		} else {
+			vgpr++;
+		}
+	} else {
+		assert(ctx->type == PIPE_SHADER_TESS_EVAL);
+		unsigned num_vgprs = uses_tes_prim_id ? 4 : 3;
+		for (unsigned i = 0; i < num_vgprs; i++) {
+			val = LLVMBuildLoad(builder, es_data[i], "");
+			ret = LLVMBuildInsertValue(builder, ret, ac_to_float(&ctx->ac, val),
+						   vgpr++, "");
+		}
+		if (num_vgprs == 3)
+			vgpr++;
+	}
+	/* Return the old thread ID. */
+	val = LLVMBuildLoad(builder, old_thread_id, "");
+	ret = LLVMBuildInsertValue(builder, ret, ac_to_float(&ctx->ac, val), vgpr++, "");
+
+	/* These two also use LDS. */
+	if (sel->info.writes_edgeflag ||
+	    (ctx->type == PIPE_SHADER_VERTEX && shader->key.mono.u.vs_export_prim_id))
+		ac_build_s_barrier(&ctx->ac);
+
+	ctx->return_value = ret;
+}
+
 /**
  * Emit the epilogue of an API VS or TES shader compiled as ESGS shader.
  */
@@ -630,7 +1258,8 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 	}
 
 	bool unterminated_es_if_block =
-		gfx10_is_ngg_passthrough(ctx->shader) &&
+		!sel->so.num_outputs &&
+		!sel->info.writes_edgeflag &&
 		!ctx->screen->use_ngg_streamout && /* no query buffer */
 		(ctx->type != PIPE_SHADER_VERTEX ||
 		 !ctx->shader->key.mono.u.vs_export_prim_id);
@@ -640,11 +1269,17 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 
 	LLVMValueRef is_gs_thread = si_is_gs_thread(ctx);
 	LLVMValueRef is_es_thread = si_is_es_thread(ctx);
-	LLVMValueRef vtxindex[] = {
-		si_unpack_param(ctx, ctx->gs_vtx01_offset, 0, 16),
-		si_unpack_param(ctx, ctx->gs_vtx01_offset, 16, 16),
-		si_unpack_param(ctx, ctx->gs_vtx23_offset, 0, 16),
-	};
+	LLVMValueRef vtxindex[3];
+
+	if (ctx->shader->key.opt.ngg_culling) {
+		vtxindex[0] = si_unpack_param(ctx, ctx->gs_vtx01_offset, 0, 9);
+		vtxindex[1] = si_unpack_param(ctx, ctx->gs_vtx01_offset, 10, 9);
+		vtxindex[2] = si_unpack_param(ctx, ctx->gs_vtx01_offset, 20, 9);
+	} else {
+		vtxindex[0] = si_unpack_param(ctx, ctx->gs_vtx01_offset, 0, 16);
+		vtxindex[1] = si_unpack_param(ctx, ctx->gs_vtx01_offset, 16, 16);
+		vtxindex[2] = si_unpack_param(ctx, ctx->gs_vtx23_offset, 0, 16);
+	}
 
 	/* Determine the number of vertices per primitive. */
 	unsigned num_vertices;
@@ -758,7 +1393,7 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 	/* Build the primitive export. */
 	if (!gfx10_ngg_export_prim_early(ctx->shader)) {
 		assert(!unterminated_es_if_block);
-		gfx10_ngg_build_export_prim(ctx, user_edgeflags);
+		gfx10_ngg_build_export_prim(ctx, user_edgeflags, NULL);
 	}
 
 	/* Export per-vertex data (positions and parameters). */
@@ -769,11 +1404,27 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 
 		/* Unconditionally (re-)load the values for proper SSA form. */
 		for (i = 0; i < info->num_outputs; i++) {
-			for (unsigned j = 0; j < 4; j++) {
-				outputs[i].values[j] =
-					LLVMBuildLoad(builder,
-						addrs[4 * i + j],
-						"");
+			/* If the NGG cull shader part computed the position, don't
+			 * use the position from the current shader part. Instead,
+			 * load it from LDS.
+			 */
+			if (info->output_semantic_name[i] == TGSI_SEMANTIC_POSITION &&
+			    ctx->shader->key.opt.ngg_culling) {
+				vertex_ptr = ngg_nogs_vertex_ptr(ctx,
+						ac_get_arg(&ctx->ac, ctx->ngg_old_thread_id));
+
+				for (unsigned j = 0; j < 4; j++) {
+					tmp = LLVMConstInt(ctx->i32, lds_pos_x + j, 0);
+					tmp = ac_build_gep0(&ctx->ac, vertex_ptr, tmp);
+					tmp = LLVMBuildLoad(builder, tmp, "");
+					outputs[i].values[j] = ac_to_float(&ctx->ac, tmp);
+				}
+			} else {
+				for (unsigned j = 0; j < 4; j++) {
+					outputs[i].values[j] =
+						LLVMBuildLoad(builder,
+							      addrs[4 * i + j], "");
+				}
 			}
 		}
 
diff --git a/src/gallium/drivers/radeonsi/si_gfx_cs.c b/src/gallium/drivers/radeonsi/si_gfx_cs.c
index 38b2abd1acb..9311b6e6386 100644
--- a/src/gallium/drivers/radeonsi/si_gfx_cs.c
+++ b/src/gallium/drivers/radeonsi/si_gfx_cs.c
@@ -445,6 +445,7 @@ void si_begin_new_gfx_cs(struct si_context *ctx)
 	ctx->last_num_tcs_input_cp = -1;
 	ctx->last_ls_hs_config = -1; /* impossible value */
 	ctx->last_binning_enabled = -1;
+	ctx->small_prim_cull_info_dirty = ctx->small_prim_cull_info_buf != NULL;
 
 	ctx->prim_discard_compute_ib_initialized = false;
 
diff --git a/src/gallium/drivers/radeonsi/si_pipe.c b/src/gallium/drivers/radeonsi/si_pipe.c
index f49961d2b43..13c13031b54 100644
--- a/src/gallium/drivers/radeonsi/si_pipe.c
+++ b/src/gallium/drivers/radeonsi/si_pipe.c
@@ -94,6 +94,8 @@ static const struct debug_named_value debug_options[] = {
 	/* 3D engine options: */
 	{ "nogfx", DBG(NO_GFX), "Disable graphics. Only multimedia compute paths can be used." },
 	{ "nongg", DBG(NO_NGG), "Disable NGG and use the legacy pipeline." },
+	{ "nggc", DBG(ALWAYS_NGG_CULLING), "Always use NGG culling even when it can hurt." },
+	{ "nonggc", DBG(NO_NGG_CULLING), "Disable NGG culling." },
 	{ "alwayspd", DBG(ALWAYS_PD), "Always enable the primitive discard compute shader." },
 	{ "pd", DBG(PD), "Enable the primitive discard compute shader for large draw calls." },
 	{ "nopd", DBG(NO_PD), "Disable the primitive discard compute shader." },
@@ -190,6 +192,7 @@ static void si_destroy_context(struct pipe_context *context)
 	si_resource_reference(&sctx->scratch_buffer, NULL);
 	si_resource_reference(&sctx->compute_scratch_buffer, NULL);
 	si_resource_reference(&sctx->wait_mem_scratch, NULL);
+	si_resource_reference(&sctx->small_prim_cull_info_buf, NULL);
 
 	si_pm4_free_state(sctx, sctx->init_config, ~0);
 	if (sctx->init_config_gs_rings)
@@ -1173,6 +1176,10 @@ radeonsi_screen_create_impl(struct radeon_winsys *ws,
 	sscreen->use_ngg = sscreen->info.chip_class >= GFX10 &&
 			   sscreen->info.family != CHIP_NAVI14 &&
 			   !(sscreen->debug_flags & DBG(NO_NGG));
+	sscreen->use_ngg_culling = sscreen->use_ngg &&
+				   !(sscreen->debug_flags & DBG(NO_NGG_CULLING));
+	sscreen->always_use_ngg_culling = sscreen->use_ngg_culling &&
+					  sscreen->debug_flags & DBG(ALWAYS_NGG_CULLING);
 	sscreen->use_ngg_streamout = false;
 
 	/* Only enable primitive binning on APUs by default. */
@@ -1305,6 +1312,7 @@ radeonsi_screen_create_impl(struct radeon_winsys *ws,
 					      4, 1, RADEON_DOMAIN_OA);
 	}
 
+	STATIC_ASSERT(sizeof(union si_vgt_stages_key) == 4);
 	return &sscreen->b;
 }
 
diff --git a/src/gallium/drivers/radeonsi/si_pipe.h b/src/gallium/drivers/radeonsi/si_pipe.h
index 563a201ec33..0a9c787dd76 100644
--- a/src/gallium/drivers/radeonsi/si_pipe.h
+++ b/src/gallium/drivers/radeonsi/si_pipe.h
@@ -183,6 +183,8 @@ enum {
 	/* 3D engine options: */
 	DBG_NO_GFX,
 	DBG_NO_NGG,
+	DBG_ALWAYS_NGG_CULLING,
+	DBG_NO_NGG_CULLING,
 	DBG_ALWAYS_PD,
 	DBG_PD,
 	DBG_NO_PD,
@@ -506,6 +508,8 @@ struct si_screen {
 	bool				dfsm_allowed;
 	bool				llvm_has_working_vgpr_indexing;
 	bool				use_ngg;
+	bool				use_ngg_culling;
+	bool				always_use_ngg_culling;
 	bool				use_ngg_streamout;
 
 	struct {
@@ -1072,6 +1076,7 @@ struct si_context {
 	bool			ls_vgpr_fix:1;
 	bool			prim_discard_cs_instancing:1;
 	bool			ngg:1;
+	uint8_t			ngg_culling;
 	int			last_index_size;
 	int			last_base_vertex;
 	int			last_start_instance;
@@ -1088,6 +1093,11 @@ struct si_context {
 	unsigned		last_vs_state;
 	enum pipe_prim_type	current_rast_prim; /* primitive type after TES, GS */
 
+	struct si_small_prim_cull_info last_small_prim_cull_info;
+	struct si_resource	*small_prim_cull_info_buf;
+	uint64_t		small_prim_cull_info_address;
+	bool			small_prim_cull_info_dirty;
+
 	/* Scratch buffer */
 	struct si_resource	*scratch_buffer;
 	unsigned		scratch_waves;
@@ -1499,6 +1509,7 @@ struct pipe_video_buffer *si_video_buffer_create(struct pipe_context *pipe,
 						 const struct pipe_video_buffer *tmpl);
 
 /* si_viewport.c */
+void si_update_ngg_small_prim_precision(struct si_context *ctx);
 void si_get_small_prim_cull_info(struct si_context *sctx,
 				 struct si_small_prim_cull_info *out);
 void si_update_vs_viewport_state(struct si_context *ctx);
diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index 24f744ba5cd..e54b9fb97ba 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -1192,7 +1192,8 @@ static void declare_vb_descriptor_input_sgprs(struct si_shader_context *ctx)
 }
 
 static void declare_vs_input_vgprs(struct si_shader_context *ctx,
-				   unsigned *num_prolog_vgprs)
+				   unsigned *num_prolog_vgprs,
+				   bool ngg_cull_shader)
 {
 	struct si_shader *shader = ctx->shader;
 
@@ -1218,6 +1219,11 @@ static void declare_vs_input_vgprs(struct si_shader_context *ctx,
 	}
 
 	if (!shader->is_gs_copy_shader) {
+		if (shader->key.opt.ngg_culling && !ngg_cull_shader) {
+			ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT,
+				   &ctx->ngg_old_thread_id);
+		}
+
 		/* Vertex load indices. */
 		if (shader->selector->info.num_inputs) {
 			ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT,
@@ -1252,12 +1258,17 @@ static void declare_vs_blit_inputs(struct si_shader_context *ctx,
 	}
 }
 
-static void declare_tes_input_vgprs(struct si_shader_context *ctx)
+static void declare_tes_input_vgprs(struct si_shader_context *ctx, bool ngg_cull_shader)
 {
 	ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->tes_u);
 	ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->tes_v);
 	ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->tes_rel_patch_id);
 	ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.tes_patch_id);
+
+	if (ctx->shader->key.opt.ngg_culling && !ngg_cull_shader) {
+		ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT,
+			   &ctx->ngg_old_thread_id);
+	}
 }
 
 enum {
@@ -1276,7 +1287,7 @@ void si_add_arg_checked(struct ac_shader_args *args,
 	ac_add_arg(args, file, registers, type, arg);
 }
 
-void si_create_function(struct si_shader_context *ctx)
+void si_create_function(struct si_shader_context *ctx, bool ngg_cull_shader)
 {
 	struct si_shader *shader = ctx->shader;
 	LLVMTypeRef returns[AC_MAX_ARGS];
@@ -1305,7 +1316,7 @@ void si_create_function(struct si_shader_context *ctx)
 			declare_vs_blit_inputs(ctx, vs_blit_property);
 
 			/* VGPRs */
-			declare_vs_input_vgprs(ctx, &num_prolog_vgprs);
+			declare_vs_input_vgprs(ctx, &num_prolog_vgprs, ngg_cull_shader);
 			break;
 		}
 
@@ -1325,7 +1336,7 @@ void si_create_function(struct si_shader_context *ctx)
 		}
 
 		/* VGPRs */
-		declare_vs_input_vgprs(ctx, &num_prolog_vgprs);
+		declare_vs_input_vgprs(ctx, &num_prolog_vgprs, ngg_cull_shader);
 
 		/* Return values */
 		if (shader->key.opt.vs_as_prim_discard_cs) {
@@ -1384,7 +1395,7 @@ void si_create_function(struct si_shader_context *ctx)
 		ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.tcs_rel_ids);
 
 		if (ctx->type == PIPE_SHADER_VERTEX) {
-			declare_vs_input_vgprs(ctx, &num_prolog_vgprs);
+			declare_vs_input_vgprs(ctx, &num_prolog_vgprs, ngg_cull_shader);
 
 			/* LS return values are inputs to the TCS main shader part. */
 			for (i = 0; i < 8 + GFX9_TCS_NUM_USER_SGPR; i++)
@@ -1419,7 +1430,8 @@ void si_create_function(struct si_shader_context *ctx)
 		ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->merged_wave_info);
 		ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_offset);
 		ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->merged_scratch_offset);
-		ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); /* unused (SPI_SHADER_PGM_LO/HI_GS << 8) */
+		ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR,
+			   &ctx->small_prim_cull_info); /* SPI_SHADER_PGM_LO_GS << 8 */
 		ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); /* unused (SPI_SHADER_PGM_LO/HI_GS >> 24) */
 
 		declare_global_desc_pointers(ctx);
@@ -1452,25 +1464,33 @@ void si_create_function(struct si_shader_context *ctx)
 		ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx45_offset);
 
 		if (ctx->type == PIPE_SHADER_VERTEX) {
-			declare_vs_input_vgprs(ctx, &num_prolog_vgprs);
+			declare_vs_input_vgprs(ctx, &num_prolog_vgprs, ngg_cull_shader);
 		} else if (ctx->type == PIPE_SHADER_TESS_EVAL) {
-			declare_tes_input_vgprs(ctx);
+			declare_tes_input_vgprs(ctx, ngg_cull_shader);
 		}
 
-		if (ctx->shader->key.as_es &&
+		if ((ctx->shader->key.as_es || ngg_cull_shader) &&
 		    (ctx->type == PIPE_SHADER_VERTEX ||
 		     ctx->type == PIPE_SHADER_TESS_EVAL)) {
-			unsigned num_user_sgprs;
+			unsigned num_user_sgprs, num_vgprs;
 
+			/* For the NGG cull shader, add 1 SGPR to hold the vertex buffer pointer. */
 			if (ctx->type == PIPE_SHADER_VERTEX)
-				num_user_sgprs = GFX9_VSGS_NUM_USER_SGPR;
+				num_user_sgprs = GFX9_VSGS_NUM_USER_SGPR + ngg_cull_shader;
 			else
 				num_user_sgprs = GFX9_TESGS_NUM_USER_SGPR;
 
+			/* The NGG cull shader has to return all 9 VGPRs + the old thread ID.
+			 *
+			 * The normal merged ESGS shader only has to return the 5 VGPRs
+			 * for the GS stage.
+			 */
+			num_vgprs = ngg_cull_shader ? 10 : 5;
+
 			/* ES return values are inputs to GS. */
 			for (i = 0; i < 8 + num_user_sgprs; i++)
 				returns[num_returns++] = ctx->i32; /* SGPRs */
-			for (i = 0; i < 5; i++)
+			for (i = 0; i < num_vgprs; i++)
 				returns[num_returns++] = ctx->f32; /* VGPRs */
 		}
 		break;
@@ -1492,7 +1512,7 @@ void si_create_function(struct si_shader_context *ctx)
 		}
 
 		/* VGPRs */
-		declare_tes_input_vgprs(ctx);
+		declare_tes_input_vgprs(ctx, ngg_cull_shader);
 		break;
 
 	case PIPE_SHADER_GEOMETRY:
@@ -1622,8 +1642,8 @@ void si_create_function(struct si_shader_context *ctx)
 		return;
 	}
 
-	si_llvm_create_func(ctx, "main", returns, num_returns,
-			    si_get_max_workgroup_size(shader));
+	si_llvm_create_func(ctx, ngg_cull_shader ? "ngg_cull_main" : "main",
+			    returns, num_returns, si_get_max_workgroup_size(shader));
 
 	/* Reserve register locations for VGPR inputs the PS prolog may need. */
 	if (ctx->type == PIPE_SHADER_FRAGMENT && !ctx->shader->is_monolithic) {
@@ -2222,6 +2242,8 @@ static void si_dump_shader_key(const struct si_shader *shader, FILE *f)
 	    !key->as_es && !key->as_ls) {
 		fprintf(f, "  opt.kill_outputs = 0x%"PRIx64"\n", key->opt.kill_outputs);
 		fprintf(f, "  opt.clip_disable = %u\n", key->opt.clip_disable);
+		if (shader_type != PIPE_SHADER_GEOMETRY)
+			fprintf(f, "  opt.ngg_culling = 0x%x\n", key->opt.ngg_culling);
 	}
 }
 
@@ -2266,7 +2288,8 @@ static bool si_vs_needs_prolog(const struct si_shader_selector *sel,
 }
 
 static bool si_build_main_function(struct si_shader_context *ctx,
-				   struct nir_shader *nir, bool free_nir)
+				   struct nir_shader *nir, bool free_nir,
+				   bool ngg_cull_shader)
 {
 	struct si_shader *shader = ctx->shader;
 	struct si_shader_selector *sel = shader->selector;
@@ -2281,6 +2304,8 @@ static bool si_build_main_function(struct si_shader_context *ctx,
 			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 if (ngg_cull_shader)
+			ctx->abi.emit_outputs = gfx10_emit_ngg_culling_epilogue_4x_wave32;
 		else if (shader->key.as_ngg)
 			ctx->abi.emit_outputs = gfx10_emit_ngg_epilogue;
 		else
@@ -2295,6 +2320,8 @@ static bool si_build_main_function(struct si_shader_context *ctx,
 
 		if (shader->key.as_es)
 			ctx->abi.emit_outputs = si_llvm_emit_es_epilogue;
+		else if (ngg_cull_shader)
+			ctx->abi.emit_outputs = gfx10_emit_ngg_culling_epilogue_4x_wave32;
 		else if (shader->key.as_ngg)
 			ctx->abi.emit_outputs = gfx10_emit_ngg_epilogue;
 		else
@@ -2314,7 +2341,7 @@ static bool si_build_main_function(struct si_shader_context *ctx,
 		return false;
 	}
 
-	si_create_function(ctx);
+	si_create_function(ctx, ngg_cull_shader);
 
 	if (ctx->shader->key.as_es || ctx->type == PIPE_SHADER_GEOMETRY)
 		si_preload_esgs_ring(ctx);
@@ -2349,6 +2376,7 @@ static bool si_build_main_function(struct si_shader_context *ctx,
 			if (sel->so.num_outputs)
 				scratch_size = 44;
 
+			assert(!ctx->gs_ngg_scratch);
 			LLVMTypeRef ai32 = LLVMArrayType(ctx->i32, scratch_size);
 			ctx->gs_ngg_scratch = LLVMAddGlobalInAddressSpace(ctx->ac.module,
 				ai32, "ngg_scratch", AC_ADDR_SPACE_LDS);
@@ -2377,7 +2405,8 @@ static bool si_build_main_function(struct si_shader_context *ctx,
 		/* This is really only needed when streamout and / or vertex
 		 * compaction is enabled.
 		 */
-		if (sel->so.num_outputs && !ctx->gs_ngg_scratch) {
+		if (!ctx->gs_ngg_scratch &&
+		    (sel->so.num_outputs || shader->key.opt.ngg_culling)) {
 			LLVMTypeRef asi32 = LLVMArrayType(ctx->i32, 8);
 			ctx->gs_ngg_scratch = LLVMAddGlobalInAddressSpace(ctx->ac.module,
 				asi32, "ngg_scratch", AC_ADDR_SPACE_LDS);
@@ -2418,19 +2447,21 @@ static bool si_build_main_function(struct si_shader_context *ctx,
 
 			if (!shader->is_monolithic ||
 			    (ctx->type == PIPE_SHADER_TESS_EVAL &&
-			     (shader->key.as_ngg && !shader->key.as_es)))
+			     shader->key.as_ngg && !shader->key.as_es &&
+			     !shader->key.opt.ngg_culling))
 				ac_init_exec_full_mask(&ctx->ac);
 
 			if ((ctx->type == PIPE_SHADER_VERTEX ||
 			     ctx->type == PIPE_SHADER_TESS_EVAL) &&
-			    shader->key.as_ngg && !shader->key.as_es) {
+			    shader->key.as_ngg && !shader->key.as_es &&
+			    !shader->key.opt.ngg_culling) {
 				gfx10_ngg_build_sendmsg_gs_alloc_req(ctx);
 
 				/* Build the primitive export at the beginning
 				 * of the shader if possible.
 				 */
 				if (gfx10_ngg_export_prim_early(shader))
-					gfx10_ngg_build_export_prim(ctx, NULL);
+					gfx10_ngg_build_export_prim(ctx, NULL, NULL);
 			}
 
 			if (ctx->type == PIPE_SHADER_TESS_CTRL ||
@@ -2500,12 +2531,14 @@ static bool si_build_main_function(struct si_shader_context *ctx,
  *
  * \param info             Shader info of the vertex shader.
  * \param num_input_sgprs  Number of input SGPRs for the vertex shader.
+ * \param has_old_  Whether the preceding shader part is the NGG cull shader.
  * \param prolog_key       Key of the VS prolog
  * \param shader_out       The vertex shader, or the next shader if merging LS+HS or ES+GS.
  * \param key              Output shader part key.
  */
 static void si_get_vs_prolog_key(const struct si_shader_info *info,
 				 unsigned num_input_sgprs,
+				 bool ngg_cull_shader,
 				 const struct si_vs_prolog_bits *prolog_key,
 				 struct si_shader *shader_out,
 				 union si_shader_part_key *key)
@@ -2518,6 +2551,9 @@ static void si_get_vs_prolog_key(const struct si_shader_info *info,
 	key->vs_prolog.as_es = shader_out->key.as_es;
 	key->vs_prolog.as_ngg = shader_out->key.as_ngg;
 
+	if (!ngg_cull_shader)
+		key->vs_prolog.has_ngg_cull_inputs = !!shader_out->key.opt.ngg_culling;
+
 	if (shader_out->selector->type == PIPE_SHADER_TESS_CTRL) {
 		key->vs_prolog.as_ls = 1;
 		key->vs_prolog.num_merged_next_stage_vgprs = 2;
@@ -2881,33 +2917,70 @@ int si_compile_shader(struct si_screen *sscreen,
 
 	shader->info.uses_instanceid = sel->info.uses_instanceid;
 
-	if (!si_build_main_function(&ctx, nir, free_nir)) {
+	LLVMValueRef ngg_cull_main_fn = NULL;
+	if (ctx.shader->key.opt.ngg_culling) {
+		if (!si_build_main_function(&ctx, nir, false, true)) {
+			si_llvm_dispose(&ctx);
+			return -1;
+		}
+		ngg_cull_main_fn = ctx.main_fn;
+		ctx.main_fn = NULL;
+		/* Re-set the IR. */
+		si_llvm_context_set_ir(&ctx, shader);
+	}
+
+	if (!si_build_main_function(&ctx, nir, free_nir, false)) {
 		si_llvm_dispose(&ctx);
 		return -1;
 	}
 
 	if (shader->is_monolithic && ctx.type == PIPE_SHADER_VERTEX) {
-		LLVMValueRef parts[2];
+		LLVMValueRef parts[4];
+		unsigned num_parts = 0;
 		bool need_prolog = si_vs_needs_prolog(sel, &shader->key.part.vs.prolog);
-
-		parts[1] = ctx.main_fn;
+		LLVMValueRef main_fn = ctx.main_fn;
+
+		if (ngg_cull_main_fn) {
+			if (need_prolog) {
+				union si_shader_part_key prolog_key;
+				si_get_vs_prolog_key(&sel->info,
+						     shader->info.num_input_sgprs,
+						     true,
+						     &shader->key.part.vs.prolog,
+						     shader, &prolog_key);
+				prolog_key.vs_prolog.is_monolithic = true;
+				si_build_vs_prolog_function(&ctx, &prolog_key);
+				parts[num_parts++] = ctx.main_fn;
+			}
+			parts[num_parts++] = ngg_cull_main_fn;
+		}
 
 		if (need_prolog) {
 			union si_shader_part_key prolog_key;
 			si_get_vs_prolog_key(&sel->info,
 					     shader->info.num_input_sgprs,
+					     false,
 					     &shader->key.part.vs.prolog,
 					     shader, &prolog_key);
 			prolog_key.vs_prolog.is_monolithic = true;
 			si_build_vs_prolog_function(&ctx, &prolog_key);
-			parts[0] = ctx.main_fn;
+			parts[num_parts++] = ctx.main_fn;
 		}
+		parts[num_parts++] = main_fn;
 
-		si_build_wrapper_function(&ctx, parts + !need_prolog,
-					  1 + need_prolog, need_prolog, 0);
+		si_build_wrapper_function(&ctx, parts, num_parts,
+					  need_prolog ? 1 : 0, 0);
 
 		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_EVAL &&
+		   ngg_cull_main_fn) {
+		LLVMValueRef parts[2];
+
+		parts[0] = ngg_cull_main_fn;
+		parts[1] = ctx.main_fn;
+
+		si_build_wrapper_function(&ctx, parts, 2, 0, 0);
 	} 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;
@@ -2935,7 +3008,7 @@ int si_compile_shader(struct si_screen *sscreen,
 			shader_ls.is_monolithic = true;
 			si_llvm_context_set_ir(&ctx, &shader_ls);
 
-			if (!si_build_main_function(&ctx, nir, free_nir)) {
+			if (!si_build_main_function(&ctx, nir, free_nir, false)) {
 				si_llvm_dispose(&ctx);
 				return -1;
 			}
@@ -2947,6 +3020,7 @@ int si_compile_shader(struct si_screen *sscreen,
 				union si_shader_part_key vs_prolog_key;
 				si_get_vs_prolog_key(&ls->info,
 						     shader_ls.info.num_input_sgprs,
+						     false,
 						     &shader->key.part.tcs.ls_prolog,
 						     shader, &vs_prolog_key);
 				vs_prolog_key.vs_prolog.is_monolithic = true;
@@ -3003,7 +3077,7 @@ int si_compile_shader(struct si_screen *sscreen,
 			shader_es.is_monolithic = true;
 			si_llvm_context_set_ir(&ctx, &shader_es);
 
-			if (!si_build_main_function(&ctx, nir, free_nir)) {
+			if (!si_build_main_function(&ctx, nir, free_nir, false)) {
 				si_llvm_dispose(&ctx);
 				return -1;
 			}
@@ -3016,6 +3090,7 @@ int si_compile_shader(struct si_screen *sscreen,
 				union si_shader_part_key vs_prolog_key;
 				si_get_vs_prolog_key(&es->info,
 						     shader_es.info.num_input_sgprs,
+						     false,
 						     &shader->key.part.gs.vs_prolog,
 						     shader, &vs_prolog_key);
 				vs_prolog_key.vs_prolog.is_monolithic = true;
@@ -3249,10 +3324,11 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx,
 	LLVMValueRef ret, func;
 	int num_returns, i;
 	unsigned first_vs_vgpr = key->vs_prolog.num_merged_next_stage_vgprs;
-	unsigned num_input_vgprs = key->vs_prolog.num_merged_next_stage_vgprs + 4;
+	unsigned num_input_vgprs = key->vs_prolog.num_merged_next_stage_vgprs + 4 +
+				   (key->vs_prolog.has_ngg_cull_inputs ? 1 : 0);
 	struct ac_arg input_sgpr_param[key->vs_prolog.num_input_sgprs];
-	struct ac_arg input_vgpr_param[9];
-	LLVMValueRef input_vgprs[9];
+	struct ac_arg input_vgpr_param[13];
+	LLVMValueRef input_vgprs[13];
 	unsigned num_all_input_regs = key->vs_prolog.num_input_sgprs +
 				      num_input_vgprs;
 	unsigned user_sgpr_base = key->vs_prolog.num_merged_next_stage_vgprs ? 8 : 0;
@@ -3427,7 +3503,7 @@ static bool si_get_vs_prolog(struct si_screen *sscreen,
 
 	/* Get the prolog. */
 	union si_shader_part_key prolog_key;
-	si_get_vs_prolog_key(&vs->info, main_part->info.num_input_sgprs,
+	si_get_vs_prolog_key(&vs->info, main_part->info.num_input_sgprs, false,
 			     key, shader, &prolog_key);
 
 	shader->prolog =
diff --git a/src/gallium/drivers/radeonsi/si_shader.h b/src/gallium/drivers/radeonsi/si_shader.h
index 295db469be3..ee1ca9cda1d 100644
--- a/src/gallium/drivers/radeonsi/si_shader.h
+++ b/src/gallium/drivers/radeonsi/si_shader.h
@@ -157,6 +157,8 @@ struct si_context;
  */
 #define SI_MAX_IO_GENERIC       32
 
+#define SI_NGG_PRIM_EDGE_FLAG_BITS ((1 << 9) | (1 << 19) | (1 << 29))
+
 /* SGPR user data indices */
 enum {
 	SI_SGPR_RW_BUFFERS,  /* rings (& stream-out, VS only) */
@@ -254,6 +256,8 @@ enum {
 #define C_VS_STATE_PROVOKING_VTX_INDEX		0xFFFFFFCF
 #define S_VS_STATE_STREAMOUT_QUERY_ENABLED(x)	(((unsigned)(x) & 0x1) << 6)
 #define C_VS_STATE_STREAMOUT_QUERY_ENABLED	0xFFFFFFBF
+#define S_VS_STATE_SMALL_PRIM_PRECISION(x)	(((unsigned)(x) & 0xF) << 7)
+#define C_VS_STATE_SMALL_PRIM_PRECISION		0xFFFFF87F
 #define S_VS_STATE_LS_OUT_PATCH_SIZE(x)		(((unsigned)(x) & 0x1FFF) << 11)
 #define C_VS_STATE_LS_OUT_PATCH_SIZE		0xFF0007FF
 #define S_VS_STATE_LS_OUT_VERTEX_SIZE(x)	(((unsigned)(x) & 0xFF) << 24)
@@ -269,6 +273,10 @@ enum {
 	SI_VS_BLIT_SGPRS_POS_TEXCOORD = 9,
 };
 
+#define SI_NGG_CULL_VIEW_SMALLPRIMS	(1 << 0) /* view.xy + small prims */
+#define SI_NGG_CULL_BACK_FACE		(1 << 1) /* back faces */
+#define SI_NGG_CULL_FRONT_FACE		(1 << 2) /* front faces */
+
 /**
  * For VS shader keys, describe any fixups required for vertex fetch.
  *
@@ -425,6 +433,7 @@ struct si_shader_selector {
 	bool		vs_needs_prolog;
 	bool		force_correct_derivs_after_kill;
 	bool		prim_discard_cs_allowed;
+	bool		ngg_culling_allowed;
 	unsigned	num_vs_inputs;
 	unsigned	num_vbos_in_user_sgprs;
 	unsigned	pa_cl_vs_out_cntl;
@@ -554,6 +563,7 @@ union si_shader_part_key {
 		unsigned	as_ls:1;
 		unsigned	as_es:1;
 		unsigned	as_ngg:1;
+		unsigned	has_ngg_cull_inputs:1; /* from the NGG cull shader */
 		/* Prologs for monolithic shaders shouldn't set EXEC. */
 		unsigned	is_monolithic:1;
 	} vs_prolog;
@@ -644,6 +654,9 @@ struct si_shader_key {
 		uint64_t	kill_outputs; /* "get_unique_index" bits */
 		unsigned	clip_disable:1;
 
+		/* For NGG VS and TES. */
+		unsigned	ngg_culling:3; /* SI_NGG_CULL_* */
+
 		/* For shaders where monolithic variants have better code.
 		 *
 		 * This is a flag that has no effect on code generation,
@@ -883,6 +896,7 @@ gfx10_is_ngg_passthrough(struct si_shader *shader)
 	return sel->type != PIPE_SHADER_GEOMETRY &&
 	       !sel->so.num_outputs &&
 	       !sel->info.writes_edgeflag &&
+	       !shader->key.opt.ngg_culling &&
 	       (sel->type != PIPE_SHADER_VERTEX ||
 		!shader->key.mono.u.vs_export_prim_id);
 }
diff --git a/src/gallium/drivers/radeonsi/si_shader_internal.h b/src/gallium/drivers/radeonsi/si_shader_internal.h
index 6509edb8181..492ebaa744b 100644
--- a/src/gallium/drivers/radeonsi/si_shader_internal.h
+++ b/src/gallium/drivers/radeonsi/si_shader_internal.h
@@ -83,6 +83,7 @@ struct si_shader_context {
 	/* Common inputs for merged shaders. */
 	struct ac_arg merged_wave_info;
 	struct ac_arg merged_scratch_offset;
+	struct ac_arg small_prim_cull_info;
 	/* API VS */
 	struct ac_arg vertex_buffers;
 	struct ac_arg vb_descriptors[5];
@@ -95,6 +96,13 @@ struct si_shader_context {
 	 *   [2:3] = NGG: output primitive type
 	 *   [4:5] = NGG: provoking vertex index
 	 *   [6]   = NGG: streamout queries enabled
+	 *   [7:10] = NGG: small prim filter precision = num_samples / quant_mode,
+	 *            but in reality it's: 1/2^n, from 1/16 to 1/4096 = 1/2^4 to 1/2^12
+	 *            Only the first 4 bits of the exponent are stored.
+	 *            Set it like this: (fui(num_samples / quant_mode) >> 23)
+	 *            Expand to FP32 like this: ((0x70 | value) << 23);
+	 *            With 0x70 = 112, we get 2^(112 + value - 127) = 2^(value - 15)
+	 *            = 1/2^(15 - value) in FP32
 	 *   [11:23] = stride between patches in DW = num_inputs * num_vertices * 4
 	 *             max = 32*32*4 + 32*4
 	 *   [24:31] = stride between vertices in DW = num_inputs * 4
@@ -102,6 +110,7 @@ struct si_shader_context {
 	 */
 	struct ac_arg vs_state_bits;
 	struct ac_arg vs_blit_inputs;
+	struct ac_arg ngg_old_thread_id; /* generated by the NGG cull shader */
 	/* HW VS */
 	struct ac_arg streamout_config;
 	struct ac_arg streamout_write_index;
@@ -297,12 +306,16 @@ void si_fix_resource_usage(struct si_screen *sscreen, struct si_shader *shader);
 void si_llvm_emit_streamout(struct si_shader_context *ctx,
 			    struct si_shader_output_values *outputs,
 			    unsigned noutput, unsigned stream);
-void si_create_function(struct si_shader_context *ctx);
+void si_create_function(struct si_shader_context *ctx, bool ngg_cull_shader);
 
 bool gfx10_ngg_export_prim_early(struct si_shader *shader);
 void gfx10_ngg_build_sendmsg_gs_alloc_req(struct si_shader_context *ctx);
 void gfx10_ngg_build_export_prim(struct si_shader_context *ctx,
-				 LLVMValueRef user_edgeflags[3]);
+				 LLVMValueRef user_edgeflags[3],
+				 LLVMValueRef prim_passthrough);
+void gfx10_emit_ngg_culling_epilogue_4x_wave32(struct ac_shader_abi *abi,
+					       unsigned max_outputs,
+					       LLVMValueRef *addrs);
 void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 			     unsigned max_outputs,
 			     LLVMValueRef *addrs);
diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm_gs.c b/src/gallium/drivers/radeonsi/si_shader_llvm_gs.c
index a622ca5b0d2..7f47f31bcdf 100644
--- a/src/gallium/drivers/radeonsi/si_shader_llvm_gs.c
+++ b/src/gallium/drivers/radeonsi/si_shader_llvm_gs.c
@@ -529,7 +529,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
 
 	builder = ctx.ac.builder;
 
-	si_create_function(&ctx);
+	si_create_function(&ctx, false);
 
 	LLVMValueRef buf_ptr = ac_get_arg(&ctx.ac, ctx.rw_buffers);
 	ctx.gsvs_ring[0] = ac_build_load_to_sgpr(&ctx.ac, buf_ptr,
diff --git a/src/gallium/drivers/radeonsi/si_state.c b/src/gallium/drivers/radeonsi/si_state.c
index c811f76e6d8..8cd56fd49d6 100644
--- a/src/gallium/drivers/radeonsi/si_state.c
+++ b/src/gallium/drivers/radeonsi/si_state.c
@@ -3101,6 +3101,7 @@ static void si_set_framebuffer_state(struct pipe_context *ctx,
 
 	si_update_ps_colorbuf0_slot(sctx);
 	si_update_poly_offset_state(sctx);
+	si_update_ngg_small_prim_precision(sctx);
 	si_mark_atom_dirty(sctx, &sctx->atoms.s.cb_render_state);
 	si_mark_atom_dirty(sctx, &sctx->atoms.s.framebuffer);
 
diff --git a/src/gallium/drivers/radeonsi/si_state_draw.c b/src/gallium/drivers/radeonsi/si_state_draw.c
index 80f5f7c943c..7f7398ff7f5 100644
--- a/src/gallium/drivers/radeonsi/si_state_draw.c
+++ b/src/gallium/drivers/radeonsi/si_state_draw.c
@@ -2038,6 +2038,45 @@ static void si_draw_vbo(struct pipe_context *ctx, const struct pipe_draw_info *i
 		sctx->do_update_shaders = true;
 	}
 
+	/* Update NGG culling settings. */
+	if (sctx->ngg &&
+	    rast_prim == PIPE_PRIM_TRIANGLES &&
+	    (sctx->screen->always_use_ngg_culling ||
+	     /* At least 1500 non-indexed triangles (4500 vertices) are needed
+	      * per draw call (no TES/GS) to enable NGG culling. Triangle strips
+	      * don't need this, because they have good reuse and therefore
+	      * perform the same as indexed triangles.
+	      */
+	     (!index_size && prim == PIPE_PRIM_TRIANGLES && direct_count > 4500 &&
+	      !sctx->tes_shader.cso && !sctx->gs_shader.cso)) &&
+	    si_get_vs(sctx)->cso->ngg_culling_allowed) {
+		unsigned ngg_culling = 0;
+
+		if (rs->rasterizer_discard) {
+			ngg_culling |= SI_NGG_CULL_FRONT_FACE |
+				       SI_NGG_CULL_BACK_FACE;
+		} else {
+			/* Polygon mode can't use view and small primitive culling,
+			 * because it draws points or lines where the culling depends
+			 * on the point or line width.
+			 */
+			if (!rs->polygon_mode_enabled)
+				ngg_culling |= SI_NGG_CULL_VIEW_SMALLPRIMS;
+
+			if (sctx->viewports.y_inverted ? rs->cull_back : rs->cull_front)
+				ngg_culling |= SI_NGG_CULL_FRONT_FACE;
+			if (sctx->viewports.y_inverted ? rs->cull_front : rs->cull_back)
+				ngg_culling |= SI_NGG_CULL_BACK_FACE;
+		}
+		if (ngg_culling != sctx->ngg_culling) {
+			sctx->ngg_culling = ngg_culling;
+			sctx->do_update_shaders = true;
+		}
+	} else if (sctx->ngg_culling) {
+		sctx->ngg_culling = false;
+		sctx->do_update_shaders = true;
+	}
+
 	if (sctx->do_update_shaders && !si_update_shaders(sctx))
 		goto return_cleanup;
 
diff --git a/src/gallium/drivers/radeonsi/si_state_shaders.c b/src/gallium/drivers/radeonsi/si_state_shaders.c
index e5ae110fa82..36dbfe9df6f 100644
--- a/src/gallium/drivers/radeonsi/si_state_shaders.c
+++ b/src/gallium/drivers/radeonsi/si_state_shaders.c
@@ -1272,8 +1272,23 @@ static void gfx10_shader_ngg(struct si_screen *sscreen, struct si_shader *shader
 	shader->ctx_reg.ngg.pa_cl_ngg_cntl =
 		S_028838_INDEX_BUF_EDGE_FLAG_ENA(gs_type == PIPE_SHADER_VERTEX);
 	shader->pa_cl_vs_out_cntl = si_get_vs_out_cntl(gs_sel, true);
+
+	/* Oversubscribe PC. This improves performance when there are too many varyings. */
+	float oversub_pc_factor = 0.25;
+
+	if (shader->key.opt.ngg_culling) {
+		/* Be more aggressive with NGG culling. */
+		if (shader->info.nr_param_exports > 4)
+			oversub_pc_factor = 1;
+		else if (shader->info.nr_param_exports > 2)
+			oversub_pc_factor = 0.75;
+		else
+			oversub_pc_factor = 0.5;
+	}
+
+	unsigned oversub_pc_lines = sscreen->info.pc_lines * oversub_pc_factor;
 	shader->ctx_reg.ngg.ge_pc_alloc = S_030980_OVERSUB_EN(1) |
-					  S_030980_NUM_PC_LINES(sscreen->info.pc_lines / 4 - 1);
+					  S_030980_NUM_PC_LINES(oversub_pc_lines - 1);
 
 	shader->ge_cntl =
 		S_03096C_PRIM_GRP_SIZE(shader->ngg.max_gsprims) |
@@ -1874,6 +1889,7 @@ static void si_shader_selector_key_hw_vs(struct si_context *sctx,
 	uint64_t linked = outputs_written & inputs_read;
 
 	key->opt.kill_outputs = ~linked & outputs_written;
+	key->opt.ngg_culling = sctx->ngg_culling;
 }
 
 /* Compute the key for the hw shader variant */
@@ -2918,6 +2934,20 @@ static void *si_create_shader_selector(struct pipe_context *ctx,
 	default:;
 	}
 
+	sel->ngg_culling_allowed =
+		sscreen->info.chip_class == GFX10 &&
+		sscreen->info.has_dedicated_vram &&
+		sscreen->use_ngg_culling &&
+		/* Disallow TES by default, because TessMark results are mixed. */
+		(sel->type == PIPE_SHADER_VERTEX ||
+		 (sscreen->always_use_ngg_culling && sel->type == PIPE_SHADER_TESS_EVAL)) &&
+		sel->info.writes_position &&
+		!sel->info.writes_viewport_index && /* cull only against viewport 0 */
+		!sel->info.writes_memory &&
+		!sel->so.num_outputs &&
+		!sel->info.properties[TGSI_PROPERTY_VS_BLIT_SGPRS_AMD] &&
+		!sel->info.properties[TGSI_PROPERTY_VS_WINDOW_SPACE_POSITION];
+
 	/* PA_CL_VS_OUT_CNTL */
 	if (sctx->chip_class <= GFX9)
 		sel->pa_cl_vs_out_cntl = si_get_vs_out_cntl(sel, false);
diff --git a/src/gallium/drivers/radeonsi/si_state_viewport.c b/src/gallium/drivers/radeonsi/si_state_viewport.c
index e0d81078139..682f00d44a8 100644
--- a/src/gallium/drivers/radeonsi/si_state_viewport.c
+++ b/src/gallium/drivers/radeonsi/si_state_viewport.c
@@ -23,10 +23,32 @@
  */
 
 #include "si_build_pm4.h"
+#include "util/u_upload_mgr.h"
 #include "util/u_viewport.h"
 
 #define SI_MAX_SCISSOR 16384
 
+void si_update_ngg_small_prim_precision(struct si_context *ctx)
+{
+	if (!ctx->screen->use_ngg_culling)
+		return;
+
+	/* Set VS_STATE.SMALL_PRIM_PRECISION for NGG culling. */
+	unsigned num_samples = ctx->framebuffer.nr_samples;
+	unsigned quant_mode = ctx->viewports.as_scissor[0].quant_mode;
+	float precision;
+
+	if (quant_mode == SI_QUANT_MODE_12_12_FIXED_POINT_1_4096TH)
+		precision = num_samples / 4096.0;
+	else if (quant_mode == SI_QUANT_MODE_14_10_FIXED_POINT_1_1024TH)
+		precision = num_samples / 1024.0;
+	else
+		precision = num_samples / 256.0;
+
+	ctx->current_vs_state &= C_VS_STATE_SMALL_PRIM_PRECISION;
+	ctx->current_vs_state |= S_VS_STATE_SMALL_PRIM_PRECISION(fui(precision) >> 23);
+}
+
 void si_get_small_prim_cull_info(struct si_context *sctx,
 				 struct si_small_prim_cull_info *out)
 {
@@ -321,6 +343,8 @@ static void si_emit_guardband(struct si_context *ctx)
 						       vp_as_scissor.quant_mode));
 	if (initial_cdw != ctx->gfx_cs->current.cdw)
 		ctx->context_roll = true;
+
+	si_update_ngg_small_prim_precision(ctx);
 }
 
 static void si_emit_scissors(struct si_context *ctx)
@@ -448,6 +472,35 @@ static void si_emit_viewports(struct si_context *ctx)
 	struct radeon_cmdbuf *cs = ctx->gfx_cs;
 	struct pipe_viewport_state *states = ctx->viewports.states;
 
+	if (ctx->screen->use_ngg_culling) {
+		/* Set the viewport info for small primitive culling. */
+		struct si_small_prim_cull_info info;
+		si_get_small_prim_cull_info(ctx, &info);
+
+		if (memcmp(&info, &ctx->last_small_prim_cull_info, sizeof(info))) {
+			unsigned offset = 0;
+
+			/* Align to 256, because the address is shifted by 8 bits. */
+			u_upload_data(ctx->b.const_uploader, 0, sizeof(info), 256,
+				      &info, &offset,
+				      (struct pipe_resource**)&ctx->small_prim_cull_info_buf);
+
+			ctx->small_prim_cull_info_address =
+				ctx->small_prim_cull_info_buf->gpu_address + offset;
+			ctx->last_small_prim_cull_info = info;
+			ctx->small_prim_cull_info_dirty = true;
+		}
+
+		if (ctx->small_prim_cull_info_dirty) {
+			/* This will end up in SGPR6 as (value << 8), shifted by the hw. */
+			radeon_add_to_buffer_list(ctx, ctx->gfx_cs, ctx->small_prim_cull_info_buf,
+						  RADEON_USAGE_READ, RADEON_PRIO_CONST_BUFFER);
+			radeon_set_sh_reg(ctx->gfx_cs, R_00B220_SPI_SHADER_PGM_LO_GS,
+					  ctx->small_prim_cull_info_address >> 8);
+			ctx->small_prim_cull_info_dirty = false;
+		}
+	}
+
 	/* The simple case: Only 1 viewport is active. */
 	if (!ctx->vs_writes_viewport_index) {
 		radeon_set_context_reg_seq(cs, R_02843C_PA_CL_VPORT_XSCALE, 6);