[Mesa-dev] [PATCH 30/92] ac/nir: start using ac_shader_abi
Bas Nieuwenhuizen
bas at basnieuwenhuizen.nl
Sun Jul 2 11:00:37 UTC 2017
Hi Nicolai,
Can we use LLVMValueRef instead of int for the shader_abi? That way we
don't force the values to be function parameters. I don't think the
shared code should have to know about that, and it is more flexible
when we want to pass those slightly differently between the two
drivers, as we can then convert them in the prologue.
- Bas
On Mon, Jun 26, 2017 at 4:10 PM, Nicolai Hähnle <nhaehnle at gmail.com> wrote:
> From: Nicolai Hähnle <nicolai.haehnle at amd.com>
>
> ---
> src/amd/common/ac_nir_to_llvm.c | 74 +++++++++++++++++++++++------------------
> 1 file changed, 41 insertions(+), 33 deletions(-)
>
> diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
> index 68865bd..e65f167 100644
> --- a/src/amd/common/ac_nir_to_llvm.c
> +++ b/src/amd/common/ac_nir_to_llvm.c
> @@ -23,20 +23,21 @@
>
> #include "ac_nir_to_llvm.h"
> #include "ac_llvm_build.h"
> #include "ac_llvm_util.h"
> #include "ac_binary.h"
> #include "sid.h"
> #include "nir/nir.h"
> #include "../vulkan/radv_descriptor_set.h"
> #include "util/bitscan.h"
> #include <llvm-c/Transforms/Scalar.h>
> +#include "ac_shader_abi.h"
> #include "ac_shader_info.h"
> #include "ac_exp_param.h"
>
> enum radeon_llvm_calling_convention {
> RADEON_LLVM_AMDGPU_VS = 87,
> RADEON_LLVM_AMDGPU_GS = 88,
> RADEON_LLVM_AMDGPU_PS = 89,
> RADEON_LLVM_AMDGPU_CS = 90,
> };
>
> @@ -50,45 +51,42 @@ enum desc_type {
> DESC_IMAGE,
> DESC_FMASK,
> DESC_SAMPLER,
> DESC_BUFFER,
> };
>
> struct nir_to_llvm_context {
> struct ac_llvm_context ac;
> const struct ac_nir_compiler_options *options;
> struct ac_shader_variant_info *shader_info;
> + struct ac_shader_abi abi;
> +
> unsigned max_workgroup_size;
> LLVMContextRef context;
> LLVMModuleRef module;
> LLVMBuilderRef builder;
> LLVMValueRef main_function;
>
> struct hash_table *defs;
> struct hash_table *phis;
>
> LLVMValueRef descriptor_sets[AC_UD_MAX_SETS];
> LLVMValueRef ring_offsets;
> LLVMValueRef push_constants;
> LLVMValueRef num_work_groups;
> LLVMValueRef workgroup_ids;
> LLVMValueRef local_invocation_ids;
> LLVMValueRef tg_size;
>
> LLVMValueRef vertex_buffers;
> - LLVMValueRef base_vertex;
> - LLVMValueRef start_instance;
> - LLVMValueRef draw_index;
> - LLVMValueRef vertex_id;
> LLVMValueRef rel_auto_id;
> LLVMValueRef vs_prim_id;
> - LLVMValueRef instance_id;
> LLVMValueRef ls_out_layout;
> LLVMValueRef es2gs_offset;
>
> LLVMValueRef tcs_offchip_layout;
> LLVMValueRef tcs_out_offsets;
> LLVMValueRef tcs_out_layout;
> LLVMValueRef tcs_in_layout;
> LLVMValueRef oc_lds;
> LLVMValueRef tess_factor_offset;
> LLVMValueRef tcs_patch_id;
> @@ -258,65 +256,70 @@ struct arg_info {
> LLVMValueRef *assign[MAX_ARGS];
> unsigned array_params_mask;
> uint8_t count;
> uint8_t user_sgpr_count;
> uint8_t sgpr_count;
> uint8_t num_user_sgprs_used;
> uint8_t num_sgprs_used;
> uint8_t num_vgprs_used;
> };
>
> -static inline void
> +static inline int
> 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;
> + int idx = info->count;
> + assert(idx < MAX_ARGS);
> + info->assign[idx] = param_ptr;
> + info->types[idx] = type;
> info->count++;
> + return idx;
> }
>
> -static inline void
> +static inline int
> add_sgpr_argument(struct arg_info *info,
> LLVMTypeRef type, LLVMValueRef *param_ptr)
> {
> - add_argument(info, type, param_ptr);
> + int idx = add_argument(info, type, param_ptr);
> info->num_sgprs_used += llvm_get_type_size(type) / 4;
> info->sgpr_count++;
> + return idx;
> }
>
> -static inline void
> +static inline int
> add_user_sgpr_argument(struct arg_info *info,
> LLVMTypeRef type,
> LLVMValueRef *param_ptr)
> {
> - add_sgpr_argument(info, type, param_ptr);
> + int idx = add_sgpr_argument(info, type, param_ptr);
> info->num_user_sgprs_used += llvm_get_type_size(type) / 4;
> info->user_sgpr_count++;
> + return idx;
> }
>
> -static inline void
> +static inline int
> add_vgpr_argument(struct arg_info *info,
> LLVMTypeRef type,
> LLVMValueRef *param_ptr)
> {
> - add_argument(info, type, param_ptr);
> + int idx = add_argument(info, type, param_ptr);
> info->num_vgprs_used += llvm_get_type_size(type) / 4;
> + return idx;
> }
>
> -static inline void
> +static inline int
> 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);
> + return 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);
> }
> @@ -739,34 +742,34 @@ static void create_function(struct nir_to_llvm_context *ctx)
> if (ctx->shader_info->info.cs.grid_components_used)
> 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)
> 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
> + ctx->abi.param_base_vertex = add_user_sgpr_argument(&args, ctx->i32, NULL); // base vertex
> + ctx->abi.param_start_instance = add_user_sgpr_argument(&args, ctx->i32, NULL);// start instance
> if (ctx->shader_info->info.vs.needs_draw_id)
> - add_user_sgpr_argument(&args, ctx->i32, &ctx->draw_index); // draw id
> + ctx->abi.param_draw_id = add_user_sgpr_argument(&args, ctx->i32, NULL); // draw id
> }
> if (ctx->options->key.vs.as_es)
> 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
> + ctx->abi.param_vertex_id = add_vgpr_argument(&args, ctx->i32, NULL); // vertex id
> if (!ctx->is_gs_copy_shader) {
> 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
> + ctx->abi.param_instance_id = add_vgpr_argument(&args, ctx->i32, NULL); // instance id
> }
> break;
> case MESA_SHADER_TESS_CTRL:
> 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
> @@ -3836,36 +3839,36 @@ static void visit_intrinsic(struct nir_to_llvm_context *ctx,
> nir_intrinsic_instr *instr)
> {
> LLVMValueRef result = NULL;
>
> switch (instr->intrinsic) {
> case nir_intrinsic_load_work_group_id: {
> result = ctx->workgroup_ids;
> break;
> }
> case nir_intrinsic_load_base_vertex: {
> - result = ctx->base_vertex;
> + result = LLVMGetParam(ctx->main_function, ctx->abi.param_base_vertex);
> break;
> }
> case nir_intrinsic_load_vertex_id_zero_base: {
> - result = ctx->vertex_id;
> + result = LLVMGetParam(ctx->main_function, ctx->abi.param_vertex_id);
> break;
> }
> case nir_intrinsic_load_local_invocation_id: {
> result = ctx->local_invocation_ids;
> break;
> }
> case nir_intrinsic_load_base_instance:
> - result = ctx->start_instance;
> + result = LLVMGetParam(ctx->main_function, ctx->abi.param_start_instance);
> break;
> case nir_intrinsic_load_draw_id:
> - result = ctx->draw_index;
> + result = LLVMGetParam(ctx->main_function, ctx->abi.param_draw_id);
> break;
> case nir_intrinsic_load_invocation_id:
> if (ctx->stage == MESA_SHADER_TESS_CTRL)
> result = unpack_param(ctx, ctx->tcs_rel_ids, 8, 5);
> else
> result = ctx->gs_invocation_id;
> break;
> case nir_intrinsic_load_primitive_id:
> if (ctx->stage == MESA_SHADER_GEOMETRY)
> result = ctx->gs_prim_id;
> @@ -3884,21 +3887,21 @@ static void visit_intrinsic(struct nir_to_llvm_context *ctx,
> ctx->shader_info->fs.force_persample = true;
> result = load_sample_pos(ctx);
> break;
> case nir_intrinsic_load_sample_mask_in:
> result = ctx->sample_coverage;
> break;
> case nir_intrinsic_load_front_face:
> result = ctx->front_face;
> break;
> case nir_intrinsic_load_instance_id:
> - result = ctx->instance_id;
> + result = LLVMGetParam(ctx->main_function, ctx->abi.param_instance_id);
> ctx->shader_info->vs.vgpr_comp_cnt = MAX2(3,
> ctx->shader_info->vs.vgpr_comp_cnt);
> break;
> case nir_intrinsic_load_num_work_groups:
> result = ctx->num_work_groups;
> break;
> case nir_intrinsic_load_local_invocation_index:
> result = visit_load_local_invocation_index(ctx);
> break;
> case nir_intrinsic_load_push_constant:
> @@ -4672,27 +4675,30 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx,
> LLVMValueRef t_list;
> LLVMValueRef input;
> LLVMValueRef buffer_index;
> int index = variable->data.location - VERT_ATTRIB_GENERIC0;
> int idx = variable->data.location;
> unsigned attrib_count = glsl_count_attribute_slots(variable->type, true);
>
> variable->data.driver_location = idx * 4;
>
> if (ctx->options->key.vs.instance_rate_inputs & (1u << index)) {
> - buffer_index = LLVMBuildAdd(ctx->builder, ctx->instance_id,
> - ctx->start_instance, "");
> + buffer_index = LLVMBuildAdd(ctx->builder,
> + LLVMGetParam(ctx->main_function, ctx->abi.param_instance_id),
> + LLVMGetParam(ctx->main_function, ctx->abi.param_start_instance), "");
> ctx->shader_info->vs.vgpr_comp_cnt = MAX2(3,
> ctx->shader_info->vs.vgpr_comp_cnt);
> - } else
> - buffer_index = LLVMBuildAdd(ctx->builder, ctx->vertex_id,
> - ctx->base_vertex, "");
> + } else {
> + buffer_index = LLVMBuildAdd(ctx->builder,
> + LLVMGetParam(ctx->main_function, ctx->abi.param_vertex_id),
> + LLVMGetParam(ctx->main_function, ctx->abi.param_base_vertex), "");
> + }
>
> for (unsigned i = 0; i < attrib_count; ++i, ++idx) {
> t_offset = LLVMConstInt(ctx->i32, index + i, false);
>
> t_list = ac_build_indexed_load_const(&ctx->ac, t_list_ptr, t_offset);
>
> input = ac_build_buffer_load_format(&ctx->ac, t_list,
> buffer_index,
> LLVMConstInt(ctx->i32, 0, false),
> true);
> @@ -6140,21 +6146,23 @@ void ac_compile_nir_shader(LLVMTargetMachineRef tm,
> default:
> break;
> }
> }
>
> static void
> ac_gs_copy_shader_emit(struct nir_to_llvm_context *ctx)
> {
> LLVMValueRef args[9];
> args[0] = ctx->gsvs_ring;
> - args[1] = LLVMBuildMul(ctx->builder, ctx->vertex_id, LLVMConstInt(ctx->i32, 4, false), "");
> + args[1] = LLVMBuildMul(ctx->builder,
> + LLVMGetParam(ctx->main_function, ctx->abi.param_vertex_id),
> + LLVMConstInt(ctx->i32, 4, false), "");
> args[3] = ctx->i32zero;
> args[4] = ctx->i32one; /* OFFEN */
> args[5] = ctx->i32zero; /* IDXEN */
> args[6] = ctx->i32one; /* GLC */
> args[7] = ctx->i32one; /* SLC */
> args[8] = ctx->i32zero; /* TFE */
>
> int idx = 0;
>
> for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
> --
> 2.9.3
>
> _______________________________________________
> mesa-dev mailing list
> mesa-dev at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
More information about the mesa-dev
mailing list