ac/nir, radv, radeonsi: Switch to using ac_shader_args

Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Acked-by: Marek Olšák <marek.olsak@amd.com>

src/amd/llvm/ac_nir_to_llvm.c

@@ -38,6 +38,7 @@
 struct ac_nir_context {
 	struct ac_llvm_context ac;
 	struct ac_shader_abi *abi;
+	const struct ac_shader_args *args;
 
 	gl_shader_stage stage;
 	shader_info *info;
@@ -1435,16 +1436,22 @@ static LLVMValueRef visit_load_push_constant(struct ac_nir_context *ctx,
 		offset += LLVMConstIntGetZExtValue(src0);
 		offset /= 4;
 
-		offset -= ctx->abi->base_inline_push_consts;
+		offset -= ctx->args->base_inline_push_consts;
 
-		if (offset + count <= ctx->abi->num_inline_push_consts) {
+		unsigned num_inline_push_consts = ctx->args->num_inline_push_consts;
+		if (offset + count <= num_inline_push_consts) {
+			LLVMValueRef push_constants[num_inline_push_consts];
+			for (unsigned i = 0; i < num_inline_push_consts; i++)
+				push_constants[i] = ac_get_arg(&ctx->ac,
+							       ctx->args->inline_push_consts[i]);
 			return ac_build_gather_values(&ctx->ac,
-						      ctx->abi->inline_push_consts + offset,
+						      push_constants + offset,
 						      count);
 		}
 	}
 
-	ptr = LLVMBuildGEP(ctx->ac.builder, ctx->abi->push_constants, &addr, 1, "");
+	ptr = LLVMBuildGEP(ctx->ac.builder,
+			   ac_get_arg(&ctx->ac, ctx->args->push_constants), &addr, 1, "");
 
 	if (instr->dest.ssa.bit_size == 8) {
 		unsigned load_dwords = instr->dest.ssa.num_components > 1 ? 2 : 1;
@@ -2902,7 +2909,8 @@ visit_load_local_invocation_index(struct ac_nir_context *ctx)
 {
 	LLVMValueRef result;
 	LLVMValueRef thread_id = ac_get_thread_id(&ctx->ac);
-	result = LLVMBuildAnd(ctx->ac.builder, ctx->abi->tg_size,
+	result = LLVMBuildAnd(ctx->ac.builder,
+			      ac_get_arg(&ctx->ac, ctx->args->tg_size),
 			      LLVMConstInt(ctx->ac.i32, 0xfc0, false), "");
 
 	if (ctx->ac.wave_size == 32)
@@ -2917,7 +2925,8 @@ visit_load_subgroup_id(struct ac_nir_context *ctx)
 {
 	if (ctx->stage == MESA_SHADER_COMPUTE) {
 		LLVMValueRef result;
-		result = LLVMBuildAnd(ctx->ac.builder, ctx->abi->tg_size,
+		result = LLVMBuildAnd(ctx->ac.builder,
+				      ac_get_arg(&ctx->ac, ctx->args->tg_size),
 				LLVMConstInt(ctx->ac.i32, 0xfc0, false), "");
 		return LLVMBuildLShr(ctx->ac.builder, result,  LLVMConstInt(ctx->ac.i32, 6, false), "");
 	} else {
@@ -2929,7 +2938,8 @@ static LLVMValueRef
 visit_load_num_subgroups(struct ac_nir_context *ctx)
 {
 	if (ctx->stage == MESA_SHADER_COMPUTE) {
-		return LLVMBuildAnd(ctx->ac.builder, ctx->abi->tg_size,
+		return LLVMBuildAnd(ctx->ac.builder,
+				    ac_get_arg(&ctx->ac, ctx->args->tg_size),
 		                    LLVMConstInt(ctx->ac.i32, 0x3f, false), "");
 	} else {
 		return LLVMConstInt(ctx->ac.i32, 1, false);
@@ -3059,8 +3069,10 @@ static LLVMValueRef load_sample_pos(struct ac_nir_context *ctx)
 	LLVMValueRef values[2];
 	LLVMValueRef pos[2];
 
-	pos[0] = ac_to_float(&ctx->ac, ctx->abi->frag_pos[0]);
+	pos[0] = ac_to_float(&ctx->ac,
-	pos[1] = ac_to_float(&ctx->ac, ctx->abi->frag_pos[1]);
+			     ac_get_arg(&ctx->ac, ctx->args->frag_pos[0]));
+	pos[1] = ac_to_float(&ctx->ac,
+			     ac_get_arg(&ctx->ac, ctx->args->frag_pos[1]));
 
 	values[0] = ac_build_fract(&ctx->ac, pos[0], 32);
 	values[1] = ac_build_fract(&ctx->ac, pos[1], 32);
@@ -3077,19 +3089,19 @@ static LLVMValueRef lookup_interp_param(struct ac_nir_context *ctx,
 	case INTERP_MODE_SMOOTH:
 	case INTERP_MODE_NONE:
 		if (location == INTERP_CENTER)
-			return ctx->abi->persp_center;
+			return ac_get_arg(&ctx->ac, ctx->args->persp_center);
 		else if (location == INTERP_CENTROID)
 			return ctx->abi->persp_centroid;
 		else if (location == INTERP_SAMPLE)
-			return ctx->abi->persp_sample;
+			return ac_get_arg(&ctx->ac, ctx->args->persp_sample);
 		break;
 	case INTERP_MODE_NOPERSPECTIVE:
 		if (location == INTERP_CENTER)
-			return ctx->abi->linear_center;
+			return ac_get_arg(&ctx->ac, ctx->args->linear_center);
 		else if (location == INTERP_CENTROID)
 			return ctx->abi->linear_centroid;
 		else if (location == INTERP_SAMPLE)
-			return ctx->abi->linear_sample;
+			return ac_get_arg(&ctx->ac, ctx->args->linear_sample);
 		break;
 	}
 	return NULL;
@@ -3203,10 +3215,10 @@ static LLVMValueRef load_interpolated_input(struct ac_nir_context *ctx,
 		LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, comp_start + comp, false);
 		if (bitsize == 16) {
 			values[comp] = ac_build_fs_interp_f16(&ctx->ac, llvm_chan, attr_number,
-							      ctx->abi->prim_mask, i, j);
+							      ac_get_arg(&ctx->ac, ctx->args->prim_mask), i, j);
 		} else {
 			values[comp] = ac_build_fs_interp(&ctx->ac, llvm_chan, attr_number,
-							  ctx->abi->prim_mask, i, j);
+							  ac_get_arg(&ctx->ac, ctx->args->prim_mask), i, j);
 		}
 	}
 
@@ -3234,7 +3246,7 @@ static LLVMValueRef load_flat_input(struct ac_nir_context *ctx,
 						      LLVMConstInt(ctx->ac.i32, 2, false),
 						      llvm_chan,
 						      attr_number,
-						      ctx->abi->prim_mask);
+						      ac_get_arg(&ctx->ac, ctx->args->prim_mask));
 		values[chan] = LLVMBuildBitCast(ctx->ac.builder, values[chan], ctx->ac.i32, "");
 		values[chan] = LLVMBuildTruncOrBitCast(ctx->ac.builder, values[chan],
 						       bit_size == 16 ? ctx->ac.i16 : ctx->ac.i32, "");
@@ -3274,8 +3286,8 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
 		LLVMValueRef values[3];
 
 		for (int i = 0; i < 3; i++) {
-			values[i] = ctx->abi->workgroup_ids[i] ?
+			values[i] = ctx->args->workgroup_ids[i].used ?
-				    ctx->abi->workgroup_ids[i] : ctx->ac.i32_0;
+				    ac_get_arg(&ctx->ac, ctx->args->workgroup_ids[i]) : ctx->ac.i32_0;
 		}
 
 		result = ac_build_gather_values(&ctx->ac, values, 3);
@@ -3289,51 +3301,56 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
 		result = ctx->abi->load_local_group_size(ctx->abi);
 		break;
 	case nir_intrinsic_load_vertex_id:
-		result = LLVMBuildAdd(ctx->ac.builder, ctx->abi->vertex_id,
+		result = LLVMBuildAdd(ctx->ac.builder,
-				      ctx->abi->base_vertex, "");
+				      ac_get_arg(&ctx->ac, ctx->args->vertex_id),
+				      ac_get_arg(&ctx->ac, ctx->args->base_vertex), "");
 		break;
 	case nir_intrinsic_load_vertex_id_zero_base: {
 		result = ctx->abi->vertex_id;
 		break;
 	}
 	case nir_intrinsic_load_local_invocation_id: {
-		result = ctx->abi->local_invocation_ids;
+		result = ac_get_arg(&ctx->ac, ctx->args->local_invocation_ids);
 		break;
 	}
 	case nir_intrinsic_load_base_instance:
-		result = ctx->abi->start_instance;
+		result = ac_get_arg(&ctx->ac, ctx->args->start_instance);
 		break;
 	case nir_intrinsic_load_draw_id:
-		result = ctx->abi->draw_id;
+		result = ac_get_arg(&ctx->ac, ctx->args->draw_id);
 		break;
 	case nir_intrinsic_load_view_index:
-		result = ctx->abi->view_index;
+		result = ac_get_arg(&ctx->ac, ctx->args->view_index);
 		break;
 	case nir_intrinsic_load_invocation_id:
 		if (ctx->stage == MESA_SHADER_TESS_CTRL) {
-			result = ac_unpack_param(&ctx->ac, ctx->abi->tcs_rel_ids, 8, 5);
+			result = ac_unpack_param(&ctx->ac,
+						 ac_get_arg(&ctx->ac, ctx->args->tcs_rel_ids),
+						 8, 5);
 		} else {
 			if (ctx->ac.chip_class >= GFX10) {
 				result = LLVMBuildAnd(ctx->ac.builder,
-						      ctx->abi->gs_invocation_id,
+						      ac_get_arg(&ctx->ac, ctx->args->gs_invocation_id),
 						      LLVMConstInt(ctx->ac.i32, 127, 0), "");
 			} else {
-				result = ctx->abi->gs_invocation_id;
+				result = ac_get_arg(&ctx->ac, ctx->args->gs_invocation_id);
 			}
 		}
 		break;
 	case nir_intrinsic_load_primitive_id:
 		if (ctx->stage == MESA_SHADER_GEOMETRY) {
-			result = ctx->abi->gs_prim_id;
+			result = ac_get_arg(&ctx->ac, ctx->args->gs_prim_id);
 		} else if (ctx->stage == MESA_SHADER_TESS_CTRL) {
-			result = ctx->abi->tcs_patch_id;
+			result = ac_get_arg(&ctx->ac, ctx->args->tcs_patch_id);
 		} else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
-			result = ctx->abi->tes_patch_id;
+			result = ac_get_arg(&ctx->ac, ctx->args->tes_patch_id);
 		} else
 			fprintf(stderr, "Unknown primitive id intrinsic: %d", ctx->stage);
 		break;
 	case nir_intrinsic_load_sample_id:
-		result = ac_unpack_param(&ctx->ac, ctx->abi->ancillary, 8, 4);
+		result = ac_unpack_param(&ctx->ac,
+					 ac_get_arg(&ctx->ac, ctx->args->ancillary),
+					 8, 4);
 		break;
 	case nir_intrinsic_load_sample_pos:
 		result = load_sample_pos(ctx);
@@ -3343,10 +3360,11 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
 		break;
 	case nir_intrinsic_load_frag_coord: {
 		LLVMValueRef values[4] = {
-			ctx->abi->frag_pos[0],
+			ac_get_arg(&ctx->ac, ctx->args->frag_pos[0]),
-			ctx->abi->frag_pos[1],
+			ac_get_arg(&ctx->ac, ctx->args->frag_pos[1]),
-			ctx->abi->frag_pos[2],
+			ac_get_arg(&ctx->ac, ctx->args->frag_pos[2]),
-			ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, ctx->abi->frag_pos[3])
+			ac_build_fdiv(&ctx->ac, ctx->ac.f32_1,
+				      ac_get_arg(&ctx->ac, ctx->args->frag_pos[3]))
 		};
 		result = ac_to_integer(&ctx->ac,
 		                       ac_build_gather_values(&ctx->ac, values, 4));
@@ -3356,7 +3374,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
 		result = ctx->abi->inputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
 		break;
 	case nir_intrinsic_load_front_face:
-		result = ctx->abi->front_face;
+		result = ac_get_arg(&ctx->ac, ctx->args->front_face);
 		break;
 	case nir_intrinsic_load_helper_invocation:
 		result = ac_build_load_helper_invocation(&ctx->ac);
@@ -3375,7 +3393,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
 		result = ctx->abi->instance_id;
 		break;
 	case nir_intrinsic_load_num_work_groups:
-		result = ctx->abi->num_work_groups;
+		result = ac_get_arg(&ctx->ac, ctx->args->num_work_groups);
 		break;
 	case nir_intrinsic_load_local_invocation_index:
 		result = visit_load_local_invocation_index(ctx);
@@ -4714,13 +4732,14 @@ setup_shared(struct ac_nir_context *ctx,
 }
 
 void ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
-		      struct nir_shader *nir)
+		      const struct ac_shader_args *args, struct nir_shader *nir)
 {
 	struct ac_nir_context ctx = {};
 	struct nir_function *func;
 
 	ctx.ac = *ac;
 	ctx.abi = abi;
+	ctx.args = args;
 
 	ctx.stage = nir->info.stage;
 	ctx.info = &nir->info;

src/amd/llvm/ac_nir_to_llvm.h

@@ -34,6 +34,7 @@ struct nir_shader;
 struct nir_variable;
 struct ac_llvm_context;
 struct ac_shader_abi;
+struct ac_shader_args;
 
 /* Interpolation locations */
 #define INTERP_CENTER 0
@@ -50,7 +51,7 @@ void ac_lower_indirect_derefs(struct nir_shader *nir, enum chip_class);
 bool ac_are_tessfactors_def_in_all_invocs(const struct nir_shader *nir);
 
 void ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
-		      struct nir_shader *nir);
+		      const struct ac_shader_args *args, struct nir_shader *nir);
 
 void
 ac_handle_shader_output_decl(struct ac_llvm_context *ctx,

src/amd/llvm/ac_shader_abi.h

@@ -25,6 +25,7 @@
 #define AC_SHADER_ABI_H
 
 #include <llvm-c/Core.h>
+#include <assert.h>
 #include "ac_shader_args.h"
 
 #include "compiler/shader_enums.h"
@@ -49,47 +50,14 @@ enum ac_descriptor_type {
  * radv to share a compiler backend.
  */
 struct ac_shader_abi {
-	LLVMValueRef base_vertex;
+	LLVMValueRef outputs[AC_LLVM_MAX_OUTPUTS * 4];
-	LLVMValueRef start_instance;
+
-	LLVMValueRef draw_id;
+	/* These input registers sometimes need to be fixed up. */
 	LLVMValueRef vertex_id;
 	LLVMValueRef instance_id;
-	LLVMValueRef tcs_patch_id;
+	LLVMValueRef persp_centroid, linear_centroid;
-	LLVMValueRef tcs_rel_ids;
+	LLVMValueRef color0, color1;
-	LLVMValueRef tes_patch_id;
-	LLVMValueRef gs_prim_id;
-	LLVMValueRef gs_invocation_id;
-
-	/* PS */
-	LLVMValueRef frag_pos[4];
-	LLVMValueRef front_face;
-	LLVMValueRef ancillary;
-	LLVMValueRef sample_coverage;
-	LLVMValueRef prim_mask;
-	LLVMValueRef color0;
-	LLVMValueRef color1;
 	LLVMValueRef user_data;
-	LLVMValueRef persp_sample;
-	LLVMValueRef persp_center;
-	LLVMValueRef persp_centroid;
-	LLVMValueRef linear_sample;
-	LLVMValueRef linear_center;
-	LLVMValueRef linear_centroid;
-
-	/* CS */
-	LLVMValueRef local_invocation_ids;
-	LLVMValueRef num_work_groups;
-	LLVMValueRef workgroup_ids[3];
-	LLVMValueRef tg_size;
-
-	/* Vulkan only */
-	LLVMValueRef push_constants;
-	LLVMValueRef inline_push_consts[AC_MAX_INLINE_PUSH_CONSTS];
-	unsigned num_inline_push_consts;
-	unsigned base_inline_push_consts;
-	LLVMValueRef view_index;
-
-	LLVMValueRef outputs[AC_LLVM_MAX_OUTPUTS * 4];
 
 	/* For VS and PS: pre-loaded shader inputs.
 	 *

src/amd/vulkan/meson.build

@@ -101,6 +101,7 @@ libradv_files = files(
   'radv_radeon_winsys.h',
   'radv_shader.c',
   'radv_shader.h',
+  'radv_shader_args.h',
   'radv_shader_helper.h',
   'radv_shader_info.c',
   'radv_query.c',

src/amd/vulkan/radv_shader_args.h

@@ -0,0 +1,76 @@
+/*
+ * Copyright © 2019 Valve Corporation.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the "Software"),
+ * to deal in the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice (including the next
+ * paragraph) shall be included in all copies or substantial portions of the
+ * Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ */
+
+#include "ac_shader_args.h"
+#include "radv_constants.h"
+#include "util/list.h"
+#include "amd_family.h"
+
+struct radv_shader_args {
+	struct ac_shader_args ac;
+	struct radv_shader_info *shader_info;
+	const struct radv_nir_compiler_options *options;
+
+	struct ac_arg descriptor_sets[MAX_SETS];
+	struct ac_arg ring_offsets;
+
+	struct ac_arg vertex_buffers;
+	struct ac_arg rel_auto_id;
+	struct ac_arg vs_prim_id;
+	struct ac_arg es2gs_offset;
+
+	struct ac_arg oc_lds;
+	struct ac_arg merged_wave_info;
+	struct ac_arg tess_factor_offset;
+	struct ac_arg tes_rel_patch_id;
+	struct ac_arg tes_u;
+	struct ac_arg tes_v;
+
+	/* HW GS */
+	/* On gfx10:
+	 *  - bits 0..10: ordered_wave_id
+	 *  - bits 12..20: number of vertices in group
+	 *  - bits 22..30: number of primitives in group
+	 */
+	struct ac_arg gs_tg_info;
+	struct ac_arg gs2vs_offset;
+	struct ac_arg gs_wave_id;
+	struct ac_arg gs_vtx_offset[6];
+
+	/* Streamout */
+	struct ac_arg streamout_buffers;
+	struct ac_arg streamout_write_idx;
+	struct ac_arg streamout_config;
+	struct ac_arg streamout_offset[4];
+
+	bool is_gs_copy_shader;
+};
+
+static inline struct radv_shader_args *
+radv_shader_args_from_ac(struct ac_shader_args *args)
+{
+	struct radv_shader_args *radv_args = NULL;
+	return (struct radv_shader_args *) container_of(args, radv_args, ac);
+}
+
+

src/gallium/drivers/radeonsi/gfx10_shader_ngg.c

@@ -31,12 +31,12 @@
 
 static LLVMValueRef get_wave_id_in_tg(struct si_shader_context *ctx)
 {
-	return si_unpack_param(ctx, ctx->param_merged_wave_info, 24, 4);
+	return si_unpack_param(ctx, ctx->merged_wave_info, 24, 4);
 }
 
 static LLVMValueRef get_tgsize(struct si_shader_context *ctx)
 {
-	return si_unpack_param(ctx, ctx->param_merged_wave_info, 28, 4);
+	return si_unpack_param(ctx, ctx->merged_wave_info, 28, 4);
 }
 
 static LLVMValueRef get_thread_id_in_tg(struct si_shader_context *ctx)
@@ -50,32 +50,22 @@ static LLVMValueRef get_thread_id_in_tg(struct si_shader_context *ctx)
 
 static LLVMValueRef ngg_get_vtx_cnt(struct si_shader_context *ctx)
 {
-	return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
+	return si_unpack_param(ctx, ctx->gs_tg_info, 12, 9);
-			    LLVMConstInt(ctx->ac.i32, 12, false),
-			    LLVMConstInt(ctx->ac.i32, 9, false),
-			    false);
 }
 
 static LLVMValueRef ngg_get_prim_cnt(struct si_shader_context *ctx)
 {
-	return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
+	return si_unpack_param(ctx, ctx->gs_tg_info, 22, 9);
-			    LLVMConstInt(ctx->ac.i32, 22, false),
-			    LLVMConstInt(ctx->ac.i32, 9, false),
-			    false);
 }
 
 static LLVMValueRef ngg_get_ordered_id(struct si_shader_context *ctx)
 {
-	return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
+	return si_unpack_param(ctx, ctx->gs_tg_info, 0, 11);
-			    ctx->i32_0,
-			    LLVMConstInt(ctx->ac.i32, 11, false),
-			    false);
 }
 
 static LLVMValueRef ngg_get_query_buf(struct si_shader_context *ctx)
 {
-	LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
+	LLVMValueRef buf_ptr = ac_get_arg(&ctx->ac, ctx->rw_buffers);
-					    ctx->param_rw_buffers);
 
 	return ac_build_load_to_sgpr(&ctx->ac, buf_ptr,
 				     LLVMConstInt(ctx->i32, GFX10_GS_QUERY_BUF, false));
@@ -212,7 +202,7 @@ static void build_streamout(struct si_shader_context *ctx,
 	struct tgsi_shader_info *info = &ctx->shader->selector->info;
 	struct pipe_stream_output_info *so = &ctx->shader->selector->so;
 	LLVMBuilderRef builder = ctx->ac.builder;
-	LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
+	LLVMValueRef buf_ptr = ac_get_arg(&ctx->ac, ctx->rw_buffers);
 	LLVMValueRef tid = get_thread_id_in_tg(ctx);
 	LLVMValueRef tmp, tmp2;
 	LLVMValueRef i32_2 = LLVMConstInt(ctx->i32, 2, false);
@@ -583,16 +573,16 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 
 	ac_build_endif(&ctx->ac, ctx->merged_wrap_if_label);
 
-	LLVMValueRef prims_in_wave = si_unpack_param(ctx, ctx->param_merged_wave_info, 8, 8);
+	LLVMValueRef prims_in_wave = si_unpack_param(ctx, ctx->merged_wave_info, 8, 8);
-	LLVMValueRef vtx_in_wave = si_unpack_param(ctx, ctx->param_merged_wave_info, 0, 8);
+	LLVMValueRef vtx_in_wave = si_unpack_param(ctx, ctx->merged_wave_info, 0, 8);
 	LLVMValueRef is_gs_thread = LLVMBuildICmp(builder, LLVMIntULT,
 						  ac_get_thread_id(&ctx->ac), prims_in_wave, "");
 	LLVMValueRef is_es_thread = LLVMBuildICmp(builder, LLVMIntULT,
 						  ac_get_thread_id(&ctx->ac), vtx_in_wave, "");
 	LLVMValueRef vtxindex[] = {
-		si_unpack_param(ctx, ctx->param_gs_vtx01_offset, 0, 16),
+		si_unpack_param(ctx, ctx->gs_vtx01_offset, 0, 16),
-		si_unpack_param(ctx, ctx->param_gs_vtx01_offset, 16, 16),
+		si_unpack_param(ctx, ctx->gs_vtx01_offset, 16, 16),
-		si_unpack_param(ctx, ctx->param_gs_vtx23_offset, 0, 16),
+		si_unpack_param(ctx, ctx->gs_vtx23_offset, 0, 16),
 	};
 
 	/* Determine the number of vertices per primitive. */
@@ -606,7 +596,7 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 			num_vertices_val = LLVMConstInt(ctx->i32, 3, 0);
 		} else {
 			/* Extract OUTPRIM field. */
-			tmp = si_unpack_param(ctx, ctx->param_vs_state_bits, 2, 2);
+			tmp = si_unpack_param(ctx, ctx->vs_state_bits, 2, 2);
 			num_vertices_val = LLVMBuildAdd(builder, tmp, ctx->i32_1, "");
 			num_vertices = 3; /* TODO: optimize for points & lines */
 		}
@@ -673,14 +663,14 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 		ac_build_ifcc(&ctx->ac, is_gs_thread, 5400);
 		/* Extract the PROVOKING_VTX_INDEX field. */
 		LLVMValueRef provoking_vtx_in_prim =
-			si_unpack_param(ctx, ctx->param_vs_state_bits, 4, 2);
+			si_unpack_param(ctx, ctx->vs_state_bits, 4, 2);
 
 		/* provoking_vtx_index = vtxindex[provoking_vtx_in_prim]; */
 		LLVMValueRef indices = ac_build_gather_values(&ctx->ac, vtxindex, 3);
 		LLVMValueRef provoking_vtx_index =
 			LLVMBuildExtractElement(builder, indices, provoking_vtx_in_prim, "");
 
-		LLVMBuildStore(builder, ctx->abi.gs_prim_id,
+		LLVMBuildStore(builder, ac_get_arg(&ctx->ac, ctx->args.gs_prim_id),
 			       ac_build_gep0(&ctx->ac, ctx->esgs_ring, provoking_vtx_index));
 		ac_build_endif(&ctx->ac, 5400);
 	}
@@ -690,7 +680,7 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 	/* Update query buffer */
 	/* TODO: this won't catch 96-bit clear_buffer via transform feedback. */
 	if (!info->properties[TGSI_PROPERTY_VS_BLIT_SGPRS_AMD]) {
-		tmp = si_unpack_param(ctx, ctx->param_vs_state_bits, 6, 1);
+		tmp = si_unpack_param(ctx, ctx->vs_state_bits, 6, 1);
 		tmp = LLVMBuildTrunc(builder, tmp, ctx->i1, "");
 		ac_build_ifcc(&ctx->ac, tmp, 5029); /* if (STREAMOUT_QUERY_ENABLED) */
 		tmp = LLVMBuildICmp(builder, LLVMIntEQ, get_wave_id_in_tg(ctx), ctx->ac.i32_0, "");
@@ -752,7 +742,8 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
 				continue;
 			}
 
-			tmp = LLVMBuildLShr(builder, ctx->abi.gs_invocation_id,
+			tmp = LLVMBuildLShr(builder,
+					    ac_get_arg(&ctx->ac, ctx->args.gs_invocation_id),
 					    LLVMConstInt(ctx->ac.i32, 8 + i, false), "");
 			prim.edgeflag[i] = LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
 
@@ -1099,7 +1090,7 @@ void gfx10_ngg_gs_emit_epilogue(struct si_shader_context *ctx)
 	}
 
 	/* Write shader query data. */
-	tmp = si_unpack_param(ctx, ctx->param_vs_state_bits, 6, 1);
+	tmp = si_unpack_param(ctx, ctx->vs_state_bits, 6, 1);
 	tmp = LLVMBuildTrunc(builder, tmp, ctx->i1, "");
 	ac_build_ifcc(&ctx->ac, tmp, 5109); /* if (STREAMOUT_QUERY_ENABLED) */
 	unsigned num_query_comps = sel->so.num_outputs ? 8 : 4;

src/gallium/drivers/radeonsi/si_compute_prim_discard.c

@@ -318,50 +318,51 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 	ac_add_function_attr(ctx->ac.context, vs, -1, AC_FUNC_ATTR_ALWAYSINLINE);
 	LLVMSetLinkage(vs, LLVMPrivateLinkage);
 
-	LLVMTypeRef const_desc_type;
+	enum ac_arg_type const_desc_type;
 	if (ctx->shader->selector->info.const_buffers_declared == 1 &&
 	    ctx->shader->selector->info.shader_buffers_declared == 0)
-		const_desc_type = ctx->f32;
+		const_desc_type = AC_ARG_CONST_FLOAT_PTR;
 	else
-		const_desc_type = ctx->v4i32;
+		const_desc_type = AC_ARG_CONST_DESC_PTR;
 
-	struct si_function_info fninfo;
+	memset(&ctx->args, 0, sizeof(ctx->args));
-	si_init_function_info(&fninfo);
 
-	LLVMValueRef index_buffers_and_constants, vertex_counter, vb_desc, const_desc;
+	struct ac_arg param_index_buffers_and_constants, param_vertex_counter;
-	LLVMValueRef base_vertex, start_instance, block_id, local_id, ordered_wave_id;
+	struct ac_arg param_vb_desc, param_const_desc;
-	LLVMValueRef restart_index, vp_scale[2], vp_translate[2], smallprim_precision;
+	struct ac_arg param_base_vertex, param_start_instance;
-	LLVMValueRef num_prims_udiv_multiplier, num_prims_udiv_terms, sampler_desc;
+	struct ac_arg param_block_id, param_local_id, param_ordered_wave_id;
-	LLVMValueRef last_wave_prim_id, vertex_count_addr;
+	struct ac_arg param_restart_index, param_smallprim_precision;
+	struct ac_arg param_num_prims_udiv_multiplier, param_num_prims_udiv_terms;
+	struct ac_arg param_sampler_desc, param_last_wave_prim_id, param_vertex_count_addr;
 
-	add_arg_assign(&fninfo, ARG_SGPR, ac_array_in_const32_addr_space(ctx->v4i32),
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR,
-		       &index_buffers_and_constants);
+		   &param_index_buffers_and_constants);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &vertex_counter);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_vertex_counter);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &last_wave_prim_id);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_last_wave_prim_id);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &vertex_count_addr);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_vertex_count_addr);
-	add_arg_assign(&fninfo, ARG_SGPR, ac_array_in_const32_addr_space(ctx->v4i32),
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR,
-		       &vb_desc);
+		   &param_vb_desc);
-	add_arg_assign(&fninfo, ARG_SGPR, ac_array_in_const32_addr_space(const_desc_type),
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, const_desc_type,
-		       &const_desc);
+		   &param_const_desc);
-	add_arg_assign(&fninfo, ARG_SGPR, ac_array_in_const32_addr_space(ctx->v8i32),
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_IMAGE_PTR,
-		       &sampler_desc);
+		   &param_sampler_desc);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &base_vertex);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_base_vertex);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &start_instance);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_start_instance);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &num_prims_udiv_multiplier);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_num_prims_udiv_multiplier);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &num_prims_udiv_terms);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_num_prims_udiv_terms);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &restart_index);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_restart_index);
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->f32, &smallprim_precision);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, &param_smallprim_precision);
 
 	/* Block ID and thread ID inputs. */
-	add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &block_id);
+	ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_block_id);
 	if (VERTEX_COUNTER_GDS_MODE == 2)
-		add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &ordered_wave_id);
+		ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_ordered_wave_id);
-	add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &local_id);
+	ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &param_local_id);
 
 	/* Create the compute shader function. */
 	unsigned old_type = ctx->type;
 	ctx->type = PIPE_SHADER_COMPUTE;
-	si_create_function(ctx, "prim_discard_cs", NULL, 0, &fninfo, THREADGROUP_SIZE);
+	si_create_function(ctx, "prim_discard_cs", NULL, 0, THREADGROUP_SIZE);
 	ctx->type = old_type;
 
 	if (VERTEX_COUNTER_GDS_MODE == 1) {
@@ -376,14 +377,14 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 
 	vs_params[num_vs_params++] = LLVMGetUndef(LLVMTypeOf(LLVMGetParam(vs, 0))); /* RW_BUFFERS */
 	vs_params[num_vs_params++] = LLVMGetUndef(LLVMTypeOf(LLVMGetParam(vs, 1))); /* BINDLESS */
-	vs_params[num_vs_params++] = const_desc;
+	vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_const_desc);
-	vs_params[num_vs_params++] = sampler_desc;
+	vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_sampler_desc);
 	vs_params[num_vs_params++] = LLVMConstInt(ctx->i32,
 					S_VS_STATE_INDEXED(key->opt.cs_indexed), 0);
-	vs_params[num_vs_params++] = base_vertex;
+	vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_base_vertex);
-	vs_params[num_vs_params++] = start_instance;
+	vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_start_instance);
 	vs_params[num_vs_params++] = ctx->i32_0; /* DrawID */
-	vs_params[num_vs_params++] = vb_desc;
+	vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_vb_desc);
 
 	vs_params[(param_vertex_id = num_vs_params++)] = NULL; /* VertexID */
 	vs_params[(param_instance_id = num_vs_params++)] = NULL; /* InstanceID */
@@ -396,6 +397,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 	/* Load descriptors. (load 8 dwords at once) */
 	LLVMValueRef input_indexbuf, output_indexbuf, tmp, desc[8];
 
+	LLVMValueRef index_buffers_and_constants = ac_get_arg(&ctx->ac, param_index_buffers_and_constants);
 	tmp = LLVMBuildPointerCast(builder, index_buffers_and_constants,
 				   ac_array_in_const32_addr_space(ctx->v8i32), "");
 	tmp = ac_build_load_to_sgpr(&ctx->ac, tmp, ctx->i32_0);
@@ -408,12 +410,17 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 
 	/* Compute PrimID and InstanceID. */
 	LLVMValueRef global_thread_id =
-		ac_build_imad(&ctx->ac, block_id,
+		ac_build_imad(&ctx->ac, ac_get_arg(&ctx->ac, param_block_id),
-			      LLVMConstInt(ctx->i32, THREADGROUP_SIZE, 0), local_id);
+			      LLVMConstInt(ctx->i32, THREADGROUP_SIZE, 0),
+			      ac_get_arg(&ctx->ac, param_local_id));
 	LLVMValueRef prim_id = global_thread_id; /* PrimID within an instance */
 	LLVMValueRef instance_id = ctx->i32_0;
 
 	if (key->opt.cs_instancing) {
+		LLVMValueRef num_prims_udiv_terms =
+			ac_get_arg(&ctx->ac, param_num_prims_udiv_terms);
+		LLVMValueRef num_prims_udiv_multiplier =
+			ac_get_arg(&ctx->ac, param_num_prims_udiv_multiplier);
 		/* Unpack num_prims_udiv_terms. */
 		LLVMValueRef post_shift = LLVMBuildAnd(builder, num_prims_udiv_terms,
 						       LLVMConstInt(ctx->i32, 0x1f, 0), "");
@@ -477,6 +484,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 		}
 	}
 
+	LLVMValueRef ordered_wave_id = ac_get_arg(&ctx->ac, param_ordered_wave_id);
+
 	/* Extract the ordered wave ID. */
 	if (VERTEX_COUNTER_GDS_MODE == 2) {
 		ordered_wave_id = LLVMBuildLShr(builder, ordered_wave_id,
@@ -485,7 +494,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 					       LLVMConstInt(ctx->i32, 0xfff, 0), "");
 	}
 	LLVMValueRef thread_id =
-		LLVMBuildAnd(builder, local_id, LLVMConstInt(ctx->i32, 63, 0), "");
+		LLVMBuildAnd(builder, ac_get_arg(&ctx->ac, param_local_id),
+			     LLVMConstInt(ctx->i32, 63, 0), "");
 
 	/* Every other triangle in a strip has a reversed vertex order, so we
 	 * need to swap vertices of odd primitives to get the correct primitive
@@ -493,6 +503,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 	 * restart complicates it, because a strip can start anywhere.
 	 */
 	LLVMValueRef prim_restart_accepted = ctx->i1true;
+	LLVMValueRef vertex_counter = ac_get_arg(&ctx->ac, param_vertex_counter);
 
 	if (key->opt.cs_prim_type == PIPE_PRIM_TRIANGLE_STRIP) {
 		/* Without primitive restart, odd primitives have reversed orientation.
@@ -520,7 +531,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 
 			for (unsigned i = 0; i < 3; i++) {
 				LLVMValueRef not_reset = LLVMBuildICmp(builder, LLVMIntNE, index[i],
-								       restart_index, "");
+								       ac_get_arg(&ctx->ac, param_restart_index),
+								       "");
 				if (i == 0)
 					index0_is_reset = LLVMBuildNot(builder, not_reset, "");
 				prim_restart_accepted = LLVMBuildAnd(builder, prim_restart_accepted,
@@ -680,6 +692,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 	LLVMValueRef vp = ac_build_load_invariant(&ctx->ac, index_buffers_and_constants,
 						  LLVMConstInt(ctx->i32, 2, 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);
@@ -699,7 +712,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 
 	LLVMValueRef accepted =
 		ac_cull_triangle(&ctx->ac, pos, prim_restart_accepted,
-				 vp_scale, vp_translate, smallprim_precision,
+				 vp_scale, vp_translate,
+				 ac_get_arg(&ctx->ac, param_smallprim_precision),
 				 &options);
 
 	LLVMValueRef accepted_threadmask = ac_get_i1_sgpr_mask(&ctx->ac, accepted);
@@ -788,7 +802,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 	if (VERTEX_COUNTER_GDS_MODE == 2) {
 		ac_build_ifcc(&ctx->ac,
 			      LLVMBuildICmp(builder, LLVMIntEQ, global_thread_id,
-					    last_wave_prim_id, ""), 12606);
+					    ac_get_arg(&ctx->ac, param_last_wave_prim_id), ""),
+			      12606);
 		LLVMValueRef count = LLVMBuildAdd(builder, start, num_prims_accepted, "");
 		count = LLVMBuildMul(builder, count,
 				     LLVMConstInt(ctx->i32, vertices_per_prim, 0), "");
@@ -798,7 +813,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 		 */
 		if (ctx->screen->info.chip_class <= GFX8) {
 			LLVMValueRef desc[] = {
-				vertex_count_addr,
+				ac_get_arg(&ctx->ac, param_vertex_count_addr),
 				LLVMConstInt(ctx->i32,
 					S_008F04_BASE_ADDRESS_HI(ctx->screen->info.address32_hi), 0),
 				LLVMConstInt(ctx->i32, 4, 0),
@@ -810,7 +825,9 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
 						    ctx->i32_0, 0, ac_glc | ac_slc, false);
 		} else {
 			LLVMBuildStore(builder, count,
-				       si_expand_32bit_pointer(ctx, vertex_count_addr));
+				       si_expand_32bit_pointer(ctx,
+							       ac_get_arg(&ctx->ac,
+									  param_vertex_count_addr)));
 		}
 		ac_build_endif(&ctx->ac, 12606);
 	} else {

src/gallium/drivers/radeonsi/si_shader_internal.h

@@ -44,22 +44,6 @@ struct pipe_debug_callback;
 #define RADEON_LLVM_MAX_SYSTEM_VALUES 11
 #define RADEON_LLVM_MAX_ADDRS 16
 
-enum si_arg_regfile {
-	ARG_SGPR,
-	ARG_VGPR
-};
-
-/**
- * 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;
-};
-
 struct si_shader_output_values {
 	LLVMValueRef values[4];
 	unsigned semantic_name;
@@ -82,6 +66,7 @@ struct si_shader_context {
 	unsigned num_images;
 	unsigned num_samplers;
 
+	struct ac_shader_args args;
 	struct ac_shader_abi abi;
 
 	/** This function is responsible for initilizing the inputs array and will be
@@ -123,19 +108,26 @@ struct si_shader_context {
 	LLVMValueRef main_fn;
 	LLVMTypeRef return_type;
 
-	/* Parameter indices for LLVMGetParam. */
+	struct ac_arg const_and_shader_buffers;
-	int param_rw_buffers;
+	struct ac_arg samplers_and_images;
-	int param_const_and_shader_buffers;
+
-	int param_samplers_and_images;
+	/* For merged shaders, the per-stage descriptors for the stage other
-	int param_bindless_samplers_and_images;
+	 * than the one we're processing, used to pass them through from the
+	 * first stage to the second.
+	 */
+	struct ac_arg other_const_and_shader_buffers;
+	struct ac_arg other_samplers_and_images;
+
+	struct ac_arg rw_buffers;
+	struct ac_arg bindless_samplers_and_images;
 	/* Common inputs for merged shaders. */
-	int param_merged_wave_info;
+	struct ac_arg merged_wave_info;
-	int param_merged_scratch_offset;
+	struct ac_arg merged_scratch_offset;
 	/* API VS */
-	int param_vertex_buffers;
+	struct ac_arg vertex_buffers;
-	int param_rel_auto_id;
+	struct ac_arg rel_auto_id;
-	int param_vs_prim_id;
+	struct ac_arg vs_prim_id;
-	int param_vertex_index0;
+	struct ac_arg vertex_index0;
 	/* VS states and layout of LS outputs / TCS inputs at the end
 	 *   [0] = clamp vertex color
 	 *   [1] = indexed
@@ -144,12 +136,12 @@ struct si_shader_context {
 	 *   [24:31] = stride between vertices in DW = num_inputs * 4
 	 *             max = 32*4
 	 */
-	int param_vs_state_bits;
+	struct ac_arg vs_state_bits;
-	int param_vs_blit_inputs;
+	struct ac_arg vs_blit_inputs;
 	/* HW VS */
-	int param_streamout_config;
+	struct ac_arg streamout_config;
-	int param_streamout_write_index;
+	struct ac_arg streamout_write_index;
-	int param_streamout_offset[4];
+	struct ac_arg streamout_offset[4];
 
 	/* API TCS & TES */
 	/* Layout of TCS outputs in the offchip buffer
@@ -161,7 +153,7 @@ struct si_shader_context {
 	 *   [12:31] = the offset of per patch attributes in the buffer in bytes.
 	 *             max = NUM_PATCHES*32*32*16
 	 */
-	int param_tcs_offchip_layout;
+	struct ac_arg tcs_offchip_layout;
 
 	/* API TCS */
 	/* Offsets where TCS outputs and TCS patch outputs live in LDS:
@@ -169,41 +161,43 @@ struct si_shader_context {
 	 *   [16:31] = TCS output patch0 offset for per-patch / 16
 	 *             max = (NUM_PATCHES + 1) * 32*32
 	 */
-	int param_tcs_out_lds_offsets;
+	struct ac_arg tcs_out_lds_offsets;
 	/* Layout of TCS outputs / TES inputs:
 	 *   [0:12] = stride between output patches in DW, num_outputs * num_vertices * 4
 	 *            max = 32*32*4 + 32*4
 	 *   [13:18] = gl_PatchVerticesIn, max = 32
 	 *   [19:31] = high 13 bits of the 32-bit address of tessellation ring buffers
 	 */
-	int param_tcs_out_lds_layout;
+	struct ac_arg tcs_out_lds_layout;
-	int param_tcs_offchip_offset;
+	struct ac_arg tcs_offchip_offset;
-	int param_tcs_factor_offset;
+	struct ac_arg tcs_factor_offset;
 
 	/* API TES */
-	int param_tes_offchip_addr;
+	struct ac_arg tes_offchip_addr;
-	int param_tes_u;
+	struct ac_arg tes_u;
-	int param_tes_v;
+	struct ac_arg tes_v;
-	int param_tes_rel_patch_id;
+	struct ac_arg tes_rel_patch_id;
 	/* HW ES */
-	int param_es2gs_offset;
+	struct ac_arg es2gs_offset;
 	/* HW GS */
 	/* On gfx10:
 	 *  - bits 0..10: ordered_wave_id
 	 *  - bits 12..20: number of vertices in group
 	 *  - bits 22..30: number of primitives in group
 	 */
-	LLVMValueRef gs_tg_info;
+	struct ac_arg gs_tg_info;
 	/* API GS */
-	int param_gs2vs_offset;
+	struct ac_arg gs2vs_offset;
-	int param_gs_wave_id; /* GFX6 */
+	struct ac_arg gs_wave_id; /* GFX6 */
-	LLVMValueRef gs_vtx_offset[6]; /* in dwords (GFX6) */
+	struct ac_arg gs_vtx_offset[6]; /* in dwords (GFX6) */
-	int param_gs_vtx01_offset; /* in dwords (GFX9) */
+	struct ac_arg gs_vtx01_offset; /* in dwords (GFX9) */
-	int param_gs_vtx23_offset; /* in dwords (GFX9) */
+	struct ac_arg gs_vtx23_offset; /* in dwords (GFX9) */
-	int param_gs_vtx45_offset; /* in dwords (GFX9) */
+	struct ac_arg gs_vtx45_offset; /* in dwords (GFX9) */
+	/* PS */
+	struct ac_arg pos_fixed_pt;
 	/* CS */
-	int param_block_size;
+	struct ac_arg block_size;
-	int param_cs_user_data;
+	struct ac_arg cs_user_data;
 
 	struct ac_llvm_compiler *compiler;
 
@@ -252,14 +246,9 @@ si_shader_context_from_abi(struct ac_shader_abi *abi)
 	return container_of(abi, ctx, abi);
 }
 
-void si_init_function_info(struct si_function_info *fninfo);
-unsigned add_arg_assign(struct si_function_info *fninfo,
-			enum si_arg_regfile regfile, LLVMTypeRef type,
-			LLVMValueRef *assign);
 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);
 unsigned si_llvm_compile(LLVMModuleRef M, struct si_shader_binary *binary,
 			 struct ac_llvm_compiler *compiler,
@@ -287,8 +276,7 @@ void si_llvm_context_set_ir(struct si_shader_context *ctx,
 
 void si_llvm_create_func(struct si_shader_context *ctx,
 			 const char *name,
-			 LLVMTypeRef *return_types, unsigned num_return_elems,
+			 LLVMTypeRef *return_types, unsigned num_return_elems);
-			 LLVMTypeRef *ParamTypes, unsigned ParamCount);
 
 void si_llvm_dispose(struct si_shader_context *ctx);
 
@@ -385,7 +373,7 @@ void si_llvm_load_input_fs(
 bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir);
 
 LLVMValueRef si_unpack_param(struct si_shader_context *ctx,
-			     unsigned param, unsigned rshift,
+			     struct ac_arg param, unsigned rshift,
 			     unsigned bitwidth);
 
 void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,

src/gallium/drivers/radeonsi/si_shader_nir.c

@@ -1047,19 +1047,19 @@ si_nir_lookup_interp_param(struct ac_shader_abi *abi,
 	case INTERP_MODE_SMOOTH:
 	case INTERP_MODE_NONE:
 		if (location == INTERP_CENTER)
-			return ctx->abi.persp_center;
+			return ac_get_arg(&ctx->ac, ctx->args.persp_center);
 		else if (location == INTERP_CENTROID)
 			return ctx->abi.persp_centroid;
 		else if (location == INTERP_SAMPLE)
-			return ctx->abi.persp_sample;
+			return ac_get_arg(&ctx->ac, ctx->args.persp_sample);
 		break;
 	case INTERP_MODE_NOPERSPECTIVE:
 		if (location == INTERP_CENTER)
-			return ctx->abi.linear_center;
+			return ac_get_arg(&ctx->ac, ctx->args.linear_center);
 		else if (location == INTERP_CENTROID)
-			return ctx->abi.linear_centroid;
+			return ac_get_arg(&ctx->ac, ctx->args.linear_centroid);
 		else if (location == INTERP_SAMPLE)
-			return ctx->abi.linear_sample;
+			return ac_get_arg(&ctx->ac, ctx->args.linear_sample);
 		break;
 	default:
 		assert(!"Unhandled interpolation mode.");
@@ -1082,8 +1082,7 @@ si_nir_load_sampler_desc(struct ac_shader_abi *abi,
 	assert(desc_type <= AC_DESC_BUFFER);
 
 	if (bindless) {
-		LLVMValueRef list =
+		LLVMValueRef list = ac_get_arg(&ctx->ac, ctx->bindless_samplers_and_images);
-			LLVMGetParam(ctx->main_fn, ctx->param_bindless_samplers_and_images);
 
 		/* dynamic_index is the bindless handle */
 		if (image) {
@@ -1114,7 +1113,7 @@ si_nir_load_sampler_desc(struct ac_shader_abi *abi,
 	unsigned num_slots = image ? ctx->num_images : ctx->num_samplers;
 	assert(const_index < num_slots || dynamic_index);
 
-	LLVMValueRef list = LLVMGetParam(ctx->main_fn, ctx->param_samplers_and_images);
+	LLVMValueRef list = ac_get_arg(&ctx->ac, ctx->samplers_and_images);
 	LLVMValueRef index = LLVMConstInt(ctx->ac.i32, const_index, false);
 
 	if (dynamic_index) {
@@ -1230,7 +1229,7 @@ bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
 			ctx->shader->key.mono.u.ps.interpolate_at_sample_force_center;
 	} else if (nir->info.stage == MESA_SHADER_COMPUTE) {
 		if (nir->info.cs.user_data_components_amd) {
-			ctx->abi.user_data = LLVMGetParam(ctx->main_fn, ctx->param_cs_user_data);
+			ctx->abi.user_data = ac_get_arg(&ctx->ac, ctx->cs_user_data);
 			ctx->abi.user_data = ac_build_expand_to_vec4(&ctx->ac, ctx->abi.user_data,
 								     nir->info.cs.user_data_components_amd);
 		}
@@ -1248,7 +1247,7 @@ bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
 		assert(gl_shader_stage_is_compute(nir->info.stage));
 		si_declare_compute_memory(ctx);
 	}
-	ac_nir_translate(&ctx->ac, &ctx->abi, nir);
+	ac_nir_translate(&ctx->ac, &ctx->abi, &ctx->args, nir);
 
 	return true;
 }

src/gallium/drivers/radeonsi/si_shader_tgsi_mem.c

@@ -224,8 +224,8 @@ image_fetch_rsrc(
 		/* Bindless descriptors are accessible from a different pair of
 		 * user SGPR indices.
 		 */
-		rsrc_ptr = LLVMGetParam(ctx->main_fn,
+		rsrc_ptr = ac_get_arg(&ctx->ac,
-					ctx->param_bindless_samplers_and_images);
+				      ctx->bindless_samplers_and_images);
 		index = lp_build_emit_fetch_src(bld_base, image, TGSI_TYPE_UNSIGNED, 0);
 
 		/* Bindless image descriptors use 16-dword slots. */
@@ -235,7 +235,7 @@ image_fetch_rsrc(
 		if (fmask)
 			index = LLVMBuildAdd(ctx->ac.builder, index, ctx->i32_1, "");
 	} else {
-		rsrc_ptr = LLVMGetParam(ctx->main_fn, ctx->param_samplers_and_images);
+		rsrc_ptr = ac_get_arg(&ctx->ac, ctx->samplers_and_images);
 
 		if (!image->Register.Indirect) {
 			index = LLVMConstInt(ctx->i32, image->Register.Index, 0);
@@ -1126,7 +1126,7 @@ static void tex_fetch_ptrs(struct lp_build_tgsi_context *bld_base,
 			   LLVMValueRef *fmask_ptr)
 {
 	struct si_shader_context *ctx = si_shader_context(bld_base);
-	LLVMValueRef list = LLVMGetParam(ctx->main_fn, ctx->param_samplers_and_images);
+	LLVMValueRef list = ac_get_arg(&ctx->ac, ctx->samplers_and_images);
 	const struct tgsi_full_instruction *inst = emit_data->inst;
 	const struct tgsi_full_src_register *reg;
 	unsigned target = inst->Texture.Texture;
@@ -1152,8 +1152,7 @@ static void tex_fetch_ptrs(struct lp_build_tgsi_context *bld_base,
 		/* Bindless descriptors are accessible from a different pair of
 		 * user SGPR indices.
 		 */
-		list = LLVMGetParam(ctx->main_fn,
+		list = ac_get_arg(&ctx->ac, ctx->bindless_samplers_and_images);
-				    ctx->param_bindless_samplers_and_images);
 		index = lp_build_emit_fetch_src(bld_base, reg,
 						TGSI_TYPE_UNSIGNED, 0);
 
@@ -1735,7 +1734,7 @@ static LLVMValueRef si_llvm_emit_fbfetch(struct si_shader_context *ctx)
 
 	/* Load the image descriptor. */
 	STATIC_ASSERT(SI_PS_IMAGE_COLORBUF0 % 2 == 0);
-	ptr = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
+	ptr = ac_get_arg(&ctx->ac, ctx->rw_buffers);
 	ptr = LLVMBuildPointerCast(ctx->ac.builder, ptr,
 				   ac_array_in_const32_addr_space(ctx->v8i32), "");
 	image = ac_build_load_to_sgpr(&ctx->ac, ptr,
@@ -1743,14 +1742,14 @@ static LLVMValueRef si_llvm_emit_fbfetch(struct si_shader_context *ctx)
 
 	unsigned chan = 0;
 
-	args.coords[chan++] = si_unpack_param(ctx, SI_PARAM_POS_FIXED_PT, 0, 16);
+	args.coords[chan++] = si_unpack_param(ctx, ctx->pos_fixed_pt, 0, 16);
 
 	if (!ctx->shader->key.mono.u.ps.fbfetch_is_1D)
-		args.coords[chan++] = si_unpack_param(ctx, SI_PARAM_POS_FIXED_PT, 16, 16);
+		args.coords[chan++] = si_unpack_param(ctx, ctx->pos_fixed_pt, 16, 16);
 
 	/* Get the current render target layer index. */
 	if (ctx->shader->key.mono.u.ps.fbfetch_layered)
-		args.coords[chan++] = si_unpack_param(ctx, SI_PARAM_ANCILLARY, 16, 11);
+		args.coords[chan++] = si_unpack_param(ctx, ctx->args.ancillary, 16, 11);
 
 	if (ctx->shader->key.mono.u.ps.fbfetch_msaa)
 		args.coords[chan++] = si_get_sample_id(ctx);

src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c

@@ -27,14 +27,6 @@
 #include "ac_llvm_util.h"
 #include "util/u_memory.h"
 
-enum si_llvm_calling_convention {
-	RADEON_LLVM_AMDGPU_VS = 87,
-	RADEON_LLVM_AMDGPU_GS = 88,
-	RADEON_LLVM_AMDGPU_PS = 89,
-	RADEON_LLVM_AMDGPU_CS = 90,
-	RADEON_LLVM_AMDGPU_HS = 93,
-};
-
 struct si_llvm_diagnostics {
 	struct pipe_debug_callback *debug;
 	unsigned retval;
@@ -1091,12 +1083,10 @@ void si_llvm_context_set_ir(struct si_shader_context *ctx,
 
 void si_llvm_create_func(struct si_shader_context *ctx,
 			 const char *name,
-			 LLVMTypeRef *return_types, unsigned num_return_elems,
+			 LLVMTypeRef *return_types, unsigned num_return_elems)
-			 LLVMTypeRef *ParamTypes, unsigned ParamCount)
 {
-	LLVMTypeRef main_fn_type, ret_type;
+	LLVMTypeRef ret_type;
-	LLVMBasicBlockRef main_fn_body;
+	enum ac_llvm_calling_convention call_conv;
-	enum si_llvm_calling_convention call_conv;
 	enum pipe_shader_type real_shader_type;
 
 	if (num_return_elems)
@@ -1106,14 +1096,6 @@ void si_llvm_create_func(struct si_shader_context *ctx,
 	else
 		ret_type = ctx->voidt;
 
-	/* Setup the function */
-	ctx->return_type = ret_type;
-	main_fn_type = LLVMFunctionType(ret_type, ParamTypes, ParamCount, 0);
-	ctx->main_fn = LLVMAddFunction(ctx->gallivm.module, name, main_fn_type);
-	main_fn_body = LLVMAppendBasicBlockInContext(ctx->ac.context,
-			ctx->main_fn, "main_body");
-	LLVMPositionBuilderAtEnd(ctx->ac.builder, main_fn_body);
-
 	real_shader_type = ctx->type;
 
 	/* LS is merged into HS (TCS), and ES is merged into GS. */
@@ -1127,25 +1109,28 @@ void si_llvm_create_func(struct si_shader_context *ctx,
 	switch (real_shader_type) {
 	case PIPE_SHADER_VERTEX:
 	case PIPE_SHADER_TESS_EVAL:
-		call_conv = RADEON_LLVM_AMDGPU_VS;
+		call_conv = AC_LLVM_AMDGPU_VS;
 		break;
 	case PIPE_SHADER_TESS_CTRL:
-		call_conv = RADEON_LLVM_AMDGPU_HS;
+		call_conv = AC_LLVM_AMDGPU_HS;
 		break;
 	case PIPE_SHADER_GEOMETRY:
-		call_conv = RADEON_LLVM_AMDGPU_GS;
+		call_conv = AC_LLVM_AMDGPU_GS;
 		break;
 	case PIPE_SHADER_FRAGMENT:
-		call_conv = RADEON_LLVM_AMDGPU_PS;
+		call_conv = AC_LLVM_AMDGPU_PS;
 		break;
 	case PIPE_SHADER_COMPUTE:
-		call_conv = RADEON_LLVM_AMDGPU_CS;
+		call_conv = AC_LLVM_AMDGPU_CS;
 		break;
 	default:
 		unreachable("Unhandle shader type");
 	}
 
-	LLVMSetFunctionCallConv(ctx->main_fn, call_conv);
+	/* Setup the function */
+	ctx->return_type = ret_type;
+	ctx->main_fn = ac_build_main(&ctx->args, &ctx->ac, call_conv, name,
+				     ret_type, ctx->gallivm.module);
 }
 
 void si_llvm_optimize_module(struct si_shader_context *ctx)