ac/nir: assign argument param pointers in one place.

Instead of having the fragile code to do a second pass, just
give the pointers you want params in to the initial code,
then call a later pass to assign them.

Reviewed-by: Bas Nieuwenhuizen <[email protected]>
Signed-off-by: Dave Airlie <[email protected]>

1 files changed, 152 insertions, 187 deletions

diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
index a939a049cae..d9bf4ea4e53 100644
--- a/src/amd/common/ac_nir_to_llvm.c
+++ b/src/amd/common/ac_nir_to_llvm.c
@@ -252,12 +252,76 @@ static void set_llvm_calling_convention(LLVMValueRef func,
 	LLVMSetFunctionCallConv(func, calling_conv);
 }
 
+#define MAX_ARGS 23
+struct arg_info {
+	LLVMTypeRef types[MAX_ARGS];
+	LLVMValueRef *assign[MAX_ARGS];
+	unsigned array_params_mask;
+	uint8_t count;
+	uint8_t user_sgpr_count;
+	uint8_t sgpr_count;
+};
+
+static inline void
+add_argument(struct arg_info *info,
+	     LLVMTypeRef type, LLVMValueRef *param_ptr)
+{
+	assert(info->count < MAX_ARGS);
+	info->assign[info->count] = param_ptr;
+	info->types[info->count] = type;
+	info->count++;
+}
+
+static inline void
+add_sgpr_argument(struct arg_info *info,
+		  LLVMTypeRef type, LLVMValueRef *param_ptr)
+{
+	add_argument(info, type, param_ptr);
+	info->sgpr_count++;
+}
+
+static inline void
+add_user_sgpr_argument(struct arg_info *info,
+		       LLVMTypeRef type,
+		       LLVMValueRef *param_ptr)
+{
+	add_sgpr_argument(info, type, param_ptr);
+	info->user_sgpr_count++;
+}
+
+static inline void
+add_vgpr_argument(struct arg_info *info,
+		  LLVMTypeRef type,
+		  LLVMValueRef *param_ptr)
+{
+	add_argument(info, type, param_ptr);
+}
+
+static inline void
+add_user_sgpr_array_argument(struct arg_info *info,
+			     LLVMTypeRef type,
+			     LLVMValueRef *param_ptr)
+{
+	info->array_params_mask |= (1 << info->count);
+	add_user_sgpr_argument(info, type, param_ptr);
+}
+
+static void assign_arguments(LLVMValueRef main_function,
+			     struct arg_info *info)
+{
+	unsigned i;
+	for (i = 0; i < info->count; i++) {
+		if (info->assign[i])
+			*info->assign[i] = LLVMGetParam(main_function, i);
+	}
+}
+
 static LLVMValueRef
 create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
                      LLVMBuilderRef builder, LLVMTypeRef *return_types,
-                     unsigned num_return_elems, LLVMTypeRef *param_types,
-                     unsigned param_count, unsigned array_params_mask,
-                     unsigned sgpr_params, unsigned max_workgroup_size,
+                     unsigned num_return_elems,
+		     struct arg_info *args,
+		     unsigned max_workgroup_size,
 		     bool unsafe_math)
 {
 	LLVMTypeRef main_function_type, ret_type;
@@ -271,7 +335,7 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
 
 	/* Setup the function */
 	main_function_type =
-	    LLVMFunctionType(ret_type, param_types, param_count, 0);
+	    LLVMFunctionType(ret_type, args->types, args->count, 0);
 	LLVMValueRef main_function =
 	    LLVMAddFunction(module, "main", main_function_type);
 	main_function_body =
@@ -279,8 +343,8 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
 	LLVMPositionBuilderAtEnd(builder, main_function_body);
 
 	LLVMSetFunctionCallConv(main_function, RADEON_LLVM_AMDGPU_CS);
-	for (unsigned i = 0; i < sgpr_params; ++i) {
-		if (array_params_mask & (1 << i)) {
+	for (unsigned i = 0; i < args->sgpr_count; ++i) {
+		if (args->array_params_mask & (1 << i)) {
 			LLVMValueRef P = LLVMGetParam(main_function, i);
 			ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_BYVAL);
 			ac_add_attr_dereferenceable(P, UINT64_MAX);
@@ -638,149 +702,128 @@ static void allocate_user_sgprs(struct nir_to_llvm_context *ctx,
 
 static void create_function(struct nir_to_llvm_context *ctx)
 {
-	LLVMTypeRef arg_types[23];
-	unsigned arg_idx = 0;
-	unsigned array_params_mask = 0;
-	unsigned sgpr_count = 0, user_sgpr_count;
 	unsigned i;
 	unsigned num_sets = ctx->options->layout ? ctx->options->layout->num_sets : 0;
 	uint8_t user_sgpr_idx;
 	struct user_sgpr_info user_sgpr_info;
+	struct arg_info args = {};
+	LLVMValueRef desc_sets;
 
 	allocate_user_sgprs(ctx, &user_sgpr_info);
 	if (user_sgpr_info.need_ring_offsets && !ctx->options->supports_spill) {
-		arg_types[arg_idx++] = const_array(ctx->v16i8, 16); /* address of rings */
+		add_user_sgpr_argument(&args, const_array(ctx->v16i8, 16), &ctx->ring_offsets); /* address of rings */
 	}
 
 	/* 1 for each descriptor set */
 	if (!user_sgpr_info.indirect_all_descriptor_sets) {
 		for (unsigned i = 0; i < num_sets; ++i) {
 			if (ctx->options->layout->set[i].layout->shader_stages & (1 << ctx->stage)) {
-				array_params_mask |= (1 << arg_idx);
-				arg_types[arg_idx++] = const_array(ctx->i8, 1024 * 1024);
+				add_user_sgpr_array_argument(&args, const_array(ctx->i8, 1024 * 1024), &ctx->descriptor_sets[i]);
 			}
 		}
-	} else {
-		array_params_mask |= (1 << arg_idx);
-		arg_types[arg_idx++] = const_array(const_array(ctx->i8, 1024 * 1024), 32);
-	}
+	} else
+		add_user_sgpr_array_argument(&args, const_array(const_array(ctx->i8, 1024 * 1024), 32), &desc_sets);
 
 	if (ctx->shader_info->info.needs_push_constants) {
 		/* 1 for push constants and dynamic descriptors */
-		array_params_mask |= (1 << arg_idx);
-		arg_types[arg_idx++] = const_array(ctx->i8, 1024 * 1024);
+		add_user_sgpr_array_argument(&args, const_array(ctx->i8, 1024 * 1024), &ctx->push_constants);
 	}
 
 	switch (ctx->stage) {
 	case MESA_SHADER_COMPUTE:
 		if (ctx->shader_info->info.cs.grid_components_used)
-			arg_types[arg_idx++] = LLVMVectorType(ctx->i32, ctx->shader_info->info.cs.grid_components_used); /* grid size */
-		user_sgpr_count = arg_idx;
-		arg_types[arg_idx++] = LLVMVectorType(ctx->i32, 3);
-		arg_types[arg_idx++] = ctx->i32;
-		sgpr_count = arg_idx;
-
-		arg_types[arg_idx++] = LLVMVectorType(ctx->i32, 3);
+			add_user_sgpr_argument(&args, LLVMVectorType(ctx->i32, ctx->shader_info->info.cs.grid_components_used), &ctx->num_work_groups); /* grid size */
+		add_sgpr_argument(&args, LLVMVectorType(ctx->i32, 3), &ctx->workgroup_ids);
+		add_sgpr_argument(&args, ctx->i32, &ctx->tg_size);
+		add_vgpr_argument(&args, LLVMVectorType(ctx->i32, 3), &ctx->local_invocation_ids);
 		break;
 	case MESA_SHADER_VERTEX:
 		if (!ctx->is_gs_copy_shader) {
 			if (ctx->shader_info->info.vs.has_vertex_buffers)
-				arg_types[arg_idx++] = const_array(ctx->v16i8, 16); /* vertex buffers */
-			arg_types[arg_idx++] = ctx->i32; // base vertex
-			arg_types[arg_idx++] = ctx->i32; // start instance
+				add_user_sgpr_argument(&args, const_array(ctx->v16i8, 16), &ctx->vertex_buffers); /* vertex buffers */
+			add_user_sgpr_argument(&args, ctx->i32, &ctx->base_vertex); // base vertex
+			add_user_sgpr_argument(&args, ctx->i32, &ctx->start_instance);// start instance
 			if (ctx->shader_info->info.vs.needs_draw_id)
-					arg_types[arg_idx++] = ctx->i32; // draw index
+				add_user_sgpr_argument(&args, ctx->i32, &ctx->draw_index); // draw id
 		}
-		user_sgpr_count = arg_idx;
 		if (ctx->options->key.vs.as_es)
-			arg_types[arg_idx++] = ctx->i32; //es2gs offset
-		else if (ctx->options->key.vs.as_ls) {
-			arg_types[arg_idx++] = ctx->i32; //ls out layout
-			user_sgpr_count++;
-		}
-		sgpr_count = arg_idx;
-		arg_types[arg_idx++] = ctx->i32; // vertex id
+			add_sgpr_argument(&args, ctx->i32, &ctx->es2gs_offset); // es2gs offset
+		else if (ctx->options->key.vs.as_ls)
+			add_user_sgpr_argument(&args, ctx->i32, &ctx->ls_out_layout); // ls out layout
+		add_vgpr_argument(&args, ctx->i32, &ctx->vertex_id); // vertex id
 		if (!ctx->is_gs_copy_shader) {
-			arg_types[arg_idx++] = ctx->i32; // rel auto id
-			arg_types[arg_idx++] = ctx->i32; // vs prim id
-			arg_types[arg_idx++] = ctx->i32; // instance id
+			add_vgpr_argument(&args, ctx->i32, &ctx->rel_auto_id); // rel auto id
+			add_vgpr_argument(&args, ctx->i32, &ctx->vs_prim_id); // vs prim id
+			add_vgpr_argument(&args, ctx->i32, &ctx->instance_id); // instance id
 		}
 		break;
 	case MESA_SHADER_TESS_CTRL:
-		arg_types[arg_idx++] = ctx->i32; // tcs offchip layout
-		arg_types[arg_idx++] = ctx->i32; // tcs out offsets
-		arg_types[arg_idx++] = ctx->i32; // tcs out layout
-		arg_types[arg_idx++] = ctx->i32; // tcs in layout
-		user_sgpr_count = arg_idx;
-		arg_types[arg_idx++] = ctx->i32; // param oc lds
-		arg_types[arg_idx++] = ctx->i32; // tess factor offset
-		sgpr_count = arg_idx;
-		arg_types[arg_idx++] = ctx->i32; // patch id
-		arg_types[arg_idx++] = ctx->i32; // rel ids;
+		add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_offchip_layout); // tcs offchip layout
+		add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_offsets); // tcs out offsets
+		add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_layout); // tcs out layout
+		add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_in_layout); // tcs in layout
+		add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // param oc lds
+		add_sgpr_argument(&args, ctx->i32, &ctx->tess_factor_offset); // tess factor offset
+		add_vgpr_argument(&args, ctx->i32, &ctx->tcs_patch_id); // patch id
+		add_vgpr_argument(&args, ctx->i32, &ctx->tcs_rel_ids); // rel ids;
 		break;
 	case MESA_SHADER_TESS_EVAL:
-		arg_types[arg_idx++] = ctx->i32; // tcs offchip layout
-		user_sgpr_count = arg_idx;
+		add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_offchip_layout); // tcs offchip layout
 		if (ctx->options->key.tes.as_es) {
-			arg_types[arg_idx++] = ctx->i32; // OC LDS
-			arg_types[arg_idx++] = ctx->i32; //
-			arg_types[arg_idx++] = ctx->i32; // es2gs offset
+			add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // OC LDS
+			add_sgpr_argument(&args, ctx->i32, NULL); //
+			add_sgpr_argument(&args, ctx->i32, &ctx->es2gs_offset); // es2gs offset
 		} else {
-			arg_types[arg_idx++] = ctx->i32; //
-			arg_types[arg_idx++] = ctx->i32; // OC LDS
+			add_sgpr_argument(&args, ctx->i32, NULL); //
+			add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // OC LDS
 		}
-		sgpr_count = arg_idx;
-		arg_types[arg_idx++] = ctx->f32; // tes_u
-		arg_types[arg_idx++] = ctx->f32; // tes_v
-		arg_types[arg_idx++] = ctx->i32; // tes rel patch id
-		arg_types[arg_idx++] = ctx->i32; // tes patch id
+		add_vgpr_argument(&args, ctx->f32, &ctx->tes_u); // tes_u
+		add_vgpr_argument(&args, ctx->f32, &ctx->tes_v); // tes_v
+		add_vgpr_argument(&args, ctx->i32, &ctx->tes_rel_patch_id); // tes rel patch id
+		add_vgpr_argument(&args, ctx->i32, &ctx->tes_patch_id); // tes patch id
 		break;
 	case MESA_SHADER_GEOMETRY:
-		arg_types[arg_idx++] = ctx->i32; // gsvs stride
-		arg_types[arg_idx++] = ctx->i32; // gsvs num entires
-		user_sgpr_count = arg_idx;
-		arg_types[arg_idx++] = ctx->i32; // gs2vs offset
-	        arg_types[arg_idx++] = ctx->i32; // wave id
-		sgpr_count = arg_idx;
-		arg_types[arg_idx++] = ctx->i32; // vtx0
-		arg_types[arg_idx++] = ctx->i32; // vtx1
-		arg_types[arg_idx++] = ctx->i32; // prim id
-		arg_types[arg_idx++] = ctx->i32; // vtx2
-		arg_types[arg_idx++] = ctx->i32; // vtx3
-		arg_types[arg_idx++] = ctx->i32; // vtx4
-		arg_types[arg_idx++] = ctx->i32; // vtx5
-		arg_types[arg_idx++] = ctx->i32; // GS instance id
+		add_user_sgpr_argument(&args, ctx->i32, &ctx->gsvs_ring_stride); // gsvs stride
+		add_user_sgpr_argument(&args, ctx->i32, &ctx->gsvs_num_entries); // gsvs num entires
+		add_sgpr_argument(&args, ctx->i32, &ctx->gs2vs_offset); // gs2vs offset
+	        add_sgpr_argument(&args, ctx->i32, &ctx->gs_wave_id); // wave id
+		add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[0]); // vtx0
+		add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[1]); // vtx1
+		add_vgpr_argument(&args, ctx->i32, &ctx->gs_prim_id); // prim id
+		add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[2]);
+		add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[3]);
+		add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[4]);
+		add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[5]);
+		add_vgpr_argument(&args, ctx->i32, &ctx->gs_invocation_id);
 		break;
 	case MESA_SHADER_FRAGMENT:
 		if (ctx->shader_info->info.ps.needs_sample_positions)
-			arg_types[arg_idx++] = ctx->i32; /* sample position offset */
-		user_sgpr_count = arg_idx;
-		arg_types[arg_idx++] = ctx->i32; /* prim mask */
-		sgpr_count = arg_idx;
-		arg_types[arg_idx++] = ctx->v2i32; /* persp sample */
-		arg_types[arg_idx++] = ctx->v2i32; /* persp center */
-		arg_types[arg_idx++] = ctx->v2i32; /* persp centroid */
-		arg_types[arg_idx++] = ctx->v3i32; /* persp pull model */
-		arg_types[arg_idx++] = ctx->v2i32; /* linear sample */
-		arg_types[arg_idx++] = ctx->v2i32; /* linear center */
-		arg_types[arg_idx++] = ctx->v2i32; /* linear centroid */
-		arg_types[arg_idx++] = ctx->f32;  /* line stipple tex */
-		arg_types[arg_idx++] = ctx->f32;  /* pos x float */
-		arg_types[arg_idx++] = ctx->f32;  /* pos y float */
-		arg_types[arg_idx++] = ctx->f32;  /* pos z float */
-		arg_types[arg_idx++] = ctx->f32;  /* pos w float */
-		arg_types[arg_idx++] = ctx->i32;  /* front face */
-		arg_types[arg_idx++] = ctx->i32;  /* ancillary */
-		arg_types[arg_idx++] = ctx->i32;  /* sample coverage */
-		arg_types[arg_idx++] = ctx->i32;  /* fixed pt */
+			add_user_sgpr_argument(&args, ctx->i32, &ctx->sample_pos_offset); /* sample position offset */
+		add_sgpr_argument(&args, ctx->i32, &ctx->prim_mask); /* prim mask */
+		add_vgpr_argument(&args, ctx->v2i32, &ctx->persp_sample); /* persp sample */
+		add_vgpr_argument(&args, ctx->v2i32, &ctx->persp_center); /* persp center */
+		add_vgpr_argument(&args, ctx->v2i32, &ctx->persp_centroid); /* persp centroid */
+		add_vgpr_argument(&args, ctx->v3i32, NULL); /* persp pull model */
+		add_vgpr_argument(&args, ctx->v2i32, &ctx->linear_sample); /* linear sample */
+		add_vgpr_argument(&args, ctx->v2i32, &ctx->linear_center); /* linear center */
+		add_vgpr_argument(&args, ctx->v2i32, &ctx->linear_centroid); /* linear centroid */
+		add_vgpr_argument(&args, ctx->f32, NULL);  /* line stipple tex */
+		add_vgpr_argument(&args, ctx->f32, &ctx->frag_pos[0]);  /* pos x float */
+		add_vgpr_argument(&args, ctx->f32, &ctx->frag_pos[1]);  /* pos y float */
+		add_vgpr_argument(&args, ctx->f32, &ctx->frag_pos[2]);  /* pos z float */
+		add_vgpr_argument(&args, ctx->f32, &ctx->frag_pos[3]);  /* pos w float */
+		add_vgpr_argument(&args, ctx->i32, &ctx->front_face);  /* front face */
+		add_vgpr_argument(&args, ctx->i32, &ctx->ancillary);  /* ancillary */
+		add_vgpr_argument(&args, ctx->i32, &ctx->sample_coverage);  /* sample coverage */
+		add_vgpr_argument(&args, ctx->i32, NULL);  /* fixed pt */
 		break;
 	default:
 		unreachable("Shader stage not implemented");
 	}
 
 	ctx->main_function = create_llvm_function(
-	    ctx->context, ctx->module, ctx->builder, NULL, 0, arg_types,
-	    arg_idx, array_params_mask, sgpr_count, ctx->max_workgroup_size,
+	    ctx->context, ctx->module, ctx->builder, NULL, 0, &args,
+	    ctx->max_workgroup_size,
 	    ctx->options->unsafe_math);
 	set_llvm_calling_convention(ctx->main_function, ctx->stage);
 
@@ -788,18 +831,19 @@ static void create_function(struct nir_to_llvm_context *ctx)
 	ctx->shader_info->num_input_vgprs = 0;
 
 	ctx->shader_info->num_user_sgprs = ctx->options->supports_spill ? 2 : 0;
-	for (i = 0; i < user_sgpr_count; i++)
-		ctx->shader_info->num_user_sgprs += llvm_get_type_size(arg_types[i]) / 4;
+	for (i = 0; i < args.user_sgpr_count; i++)
+		ctx->shader_info->num_user_sgprs += llvm_get_type_size(args.types[i]) / 4;
 
 	ctx->shader_info->num_input_sgprs = ctx->shader_info->num_user_sgprs;
-	for (; i < sgpr_count; i++)
-		ctx->shader_info->num_input_sgprs += llvm_get_type_size(arg_types[i]) / 4;
+	for (; i < args.sgpr_count; i++)
+		ctx->shader_info->num_input_sgprs += llvm_get_type_size(args.types[i]) / 4;
 
 	if (ctx->stage != MESA_SHADER_FRAGMENT)
-		for (; i < arg_idx; ++i)
-			ctx->shader_info->num_input_vgprs += llvm_get_type_size(arg_types[i]) / 4;
+		for (; i < args.count; ++i)
+			ctx->shader_info->num_input_vgprs += llvm_get_type_size(args.types[i]) / 4;
+
+	assign_arguments(ctx->main_function, &args);
 
-	arg_idx = 0;
 	user_sgpr_idx = 0;
 
 	if (ctx->options->supports_spill || user_sgpr_info.need_ring_offsets) {
@@ -810,22 +854,18 @@ static void create_function(struct nir_to_llvm_context *ctx)
 							       NULL, 0, AC_FUNC_ATTR_READNONE);
 			ctx->ring_offsets = LLVMBuildBitCast(ctx->builder, ctx->ring_offsets,
 							     const_array(ctx->v16i8, 16), "");
-		} else
-			ctx->ring_offsets = LLVMGetParam(ctx->main_function, arg_idx++);
+		}
 	}
 
 	if (!user_sgpr_info.indirect_all_descriptor_sets) {
 		for (unsigned i = 0; i < num_sets; ++i) {
 			if (ctx->options->layout->set[i].layout->shader_stages & (1 << ctx->stage)) {
 				set_userdata_location(&ctx->shader_info->user_sgprs_locs.descriptor_sets[i], &user_sgpr_idx, 2);
-				ctx->descriptor_sets[i] =
-					LLVMGetParam(ctx->main_function, arg_idx++);
 			} else
 				ctx->descriptor_sets[i] = NULL;
 		}
 	} else {
 		uint32_t desc_sgpr_idx = user_sgpr_idx;
-		LLVMValueRef desc_sets = LLVMGetParam(ctx->main_function, arg_idx++);
 		set_userdata_location_shader(ctx, AC_UD_INDIRECT_DESCRIPTOR_SETS, &user_sgpr_idx, 2);
 
 		for (unsigned i = 0; i < num_sets; ++i) {
@@ -840,7 +880,6 @@ static void create_function(struct nir_to_llvm_context *ctx)
 	}
 
 	if (ctx->shader_info->info.needs_push_constants) {
-		ctx->push_constants = LLVMGetParam(ctx->main_function, arg_idx++);
 		set_userdata_location_shader(ctx, AC_UD_PUSH_CONSTANTS, &user_sgpr_idx, 2);
 	}
 
@@ -848,113 +887,39 @@ static void create_function(struct nir_to_llvm_context *ctx)
 	case MESA_SHADER_COMPUTE:
 		if (ctx->shader_info->info.cs.grid_components_used) {
 			set_userdata_location_shader(ctx, AC_UD_CS_GRID_SIZE, &user_sgpr_idx, ctx->shader_info->info.cs.grid_components_used);
-			ctx->num_work_groups =
-				LLVMGetParam(ctx->main_function, arg_idx++);
 		}
-		ctx->workgroup_ids =
-		    LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tg_size =
-		    LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->local_invocation_ids =
-		    LLVMGetParam(ctx->main_function, arg_idx++);
 		break;
 	case MESA_SHADER_VERTEX:
 		if (!ctx->is_gs_copy_shader) {
 			if (ctx->shader_info->info.vs.has_vertex_buffers) {
 				set_userdata_location_shader(ctx, AC_UD_VS_VERTEX_BUFFERS, &user_sgpr_idx, 2);
-				ctx->vertex_buffers = LLVMGetParam(ctx->main_function, arg_idx++);
 			}
 			unsigned vs_num = 2;
 			if (ctx->shader_info->info.vs.needs_draw_id)
 				vs_num++;
 
 			set_userdata_location_shader(ctx, AC_UD_VS_BASE_VERTEX_START_INSTANCE, &user_sgpr_idx, vs_num);
-
-			ctx->base_vertex = LLVMGetParam(ctx->main_function, arg_idx++);
-			ctx->start_instance = LLVMGetParam(ctx->main_function, arg_idx++);
-			if (ctx->shader_info->info.vs.needs_draw_id)
-				ctx->draw_index = LLVMGetParam(ctx->main_function, arg_idx++);
 		}
-		if (ctx->options->key.vs.as_es)
-			ctx->es2gs_offset = LLVMGetParam(ctx->main_function, arg_idx++);
-		else if (ctx->options->key.vs.as_ls) {
+		if (ctx->options->key.vs.as_ls) {
 			set_userdata_location_shader(ctx, AC_UD_VS_LS_TCS_IN_LAYOUT, &user_sgpr_idx, 1);
-			ctx->ls_out_layout = LLVMGetParam(ctx->main_function, arg_idx++);
-		}
-		ctx->vertex_id = LLVMGetParam(ctx->main_function, arg_idx++);
-		if (!ctx->is_gs_copy_shader) {
-			ctx->rel_auto_id = LLVMGetParam(ctx->main_function, arg_idx++);
-			ctx->vs_prim_id = LLVMGetParam(ctx->main_function, arg_idx++);
-			ctx->instance_id = LLVMGetParam(ctx->main_function, arg_idx++);
 		}
 		if (ctx->options->key.vs.as_ls)
 			declare_tess_lds(ctx);
 		break;
 	case MESA_SHADER_TESS_CTRL:
 		set_userdata_location_shader(ctx, AC_UD_TCS_OFFCHIP_LAYOUT, &user_sgpr_idx, 4);
-		ctx->tcs_offchip_layout = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tcs_out_offsets = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tcs_out_layout = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tcs_in_layout = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->oc_lds = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tess_factor_offset = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tcs_patch_id = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tcs_rel_ids = LLVMGetParam(ctx->main_function, arg_idx++);
-
 		declare_tess_lds(ctx);
 		break;
 	case MESA_SHADER_TESS_EVAL:
 		set_userdata_location_shader(ctx, AC_UD_TES_OFFCHIP_LAYOUT, &user_sgpr_idx, 1);
-		ctx->tcs_offchip_layout = LLVMGetParam(ctx->main_function, arg_idx++);
-		if (ctx->options->key.tes.as_es) {
-			ctx->oc_lds = LLVMGetParam(ctx->main_function, arg_idx++);
-			arg_idx++;
-			ctx->es2gs_offset = LLVMGetParam(ctx->main_function, arg_idx++);
-		} else {
-			arg_idx++;
-			ctx->oc_lds = LLVMGetParam(ctx->main_function, arg_idx++);
-		}
-		ctx->tes_u = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tes_v = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tes_rel_patch_id = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->tes_patch_id = LLVMGetParam(ctx->main_function, arg_idx++);
 		break;
 	case MESA_SHADER_GEOMETRY:
 		set_userdata_location_shader(ctx, AC_UD_GS_VS_RING_STRIDE_ENTRIES, &user_sgpr_idx, 2);
-		ctx->gsvs_ring_stride = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gsvs_num_entries = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs2vs_offset = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_wave_id = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_vtx_offset[0] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_vtx_offset[1] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_prim_id = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_vtx_offset[2] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_vtx_offset[3] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_vtx_offset[4] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_vtx_offset[5] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->gs_invocation_id = LLVMGetParam(ctx->main_function, arg_idx++);
 		break;
 	case MESA_SHADER_FRAGMENT:
 		if (ctx->shader_info->info.ps.needs_sample_positions) {
 			set_userdata_location_shader(ctx, AC_UD_PS_SAMPLE_POS_OFFSET, &user_sgpr_idx, 1);
-			ctx->sample_pos_offset = LLVMGetParam(ctx->main_function, arg_idx++);
 		}
-		ctx->prim_mask = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->persp_sample = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->persp_center = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->persp_centroid = LLVMGetParam(ctx->main_function, arg_idx++);
-		arg_idx++;
-		ctx->linear_sample = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->linear_center = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->linear_centroid = LLVMGetParam(ctx->main_function, arg_idx++);
-		arg_idx++; /* line stipple */
-		ctx->frag_pos[0] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->frag_pos[1] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->frag_pos[2] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->frag_pos[3] = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->front_face = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->ancillary = LLVMGetParam(ctx->main_function, arg_idx++);
-		ctx->sample_coverage = LLVMGetParam(ctx->main_function, arg_idx++);
 		break;
 	default:
 		unreachable("Shader stage not implemented");