On 2 November 2017 at 11:50, Timothy Arceri <tarc...@itsqueeze.com> wrote:
> ---
Yeah why not,
For the series:
Reviewed-by: Dave Airlie <airl...@redhat.com>
> src/amd/common/ac_nir_to_llvm.c | 360
> ++++++++++++++++++++--------------------
> 1 file changed, 179 insertions(+), 181 deletions(-)
>
> diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
> index 2ec30517e0..d792042925 100644
> --- a/src/amd/common/ac_nir_to_llvm.c
> +++ b/src/amd/common/ac_nir_to_llvm.c
> @@ -130,21 +130,20 @@ struct nir_to_llvm_context {
> LLVMValueRef hs_ring_tess_factor;
>
> LLVMValueRef prim_mask;
> LLVMValueRef sample_pos_offset;
> LLVMValueRef persp_sample, persp_center, persp_centroid;
> LLVMValueRef linear_sample, linear_center, linear_centroid;
>
> LLVMTypeRef i1;
> LLVMTypeRef i8;
> LLVMTypeRef i16;
> - LLVMTypeRef i32;
> LLVMTypeRef i64;
> LLVMTypeRef v2i32;
> LLVMTypeRef v3i32;
> LLVMTypeRef v4i32;
> LLVMTypeRef v8i32;
> LLVMTypeRef f64;
> LLVMTypeRef f32;
> LLVMTypeRef f16;
> LLVMTypeRef v2f32;
> LLVMTypeRef v4f32;
> @@ -461,29 +460,29 @@ static LLVMValueRef
> get_tcs_out_patch_stride(struct nir_to_llvm_context *ctx)
> {
> return unpack_param(&ctx->ac, ctx->tcs_out_layout, 0, 13);
> }
>
> static LLVMValueRef
> get_tcs_out_patch0_offset(struct nir_to_llvm_context *ctx)
> {
> return LLVMBuildMul(ctx->builder,
> unpack_param(&ctx->ac, ctx->tcs_out_offsets, 0,
> 16),
> - LLVMConstInt(ctx->i32, 4, false), "");
> + LLVMConstInt(ctx->ac.i32, 4, false), "");
> }
>
> static LLVMValueRef
> get_tcs_out_patch0_patch_data_offset(struct nir_to_llvm_context *ctx)
> {
> return LLVMBuildMul(ctx->builder,
> unpack_param(&ctx->ac, ctx->tcs_out_offsets, 16,
> 16),
> - LLVMConstInt(ctx->i32, 4, false), "");
> + LLVMConstInt(ctx->ac.i32, 4, false), "");
> }
>
> static LLVMValueRef
> get_tcs_in_current_patch_offset(struct nir_to_llvm_context *ctx)
> {
> LLVMValueRef patch_stride = get_tcs_in_patch_stride(ctx);
> LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
>
> return LLVMBuildMul(ctx->builder, patch_stride, rel_patch_id, "");
> }
> @@ -663,21 +662,21 @@ radv_define_common_user_sgprs_phase2(struct
> nir_to_llvm_context *ctx,
> } else
> ctx->descriptor_sets[i] = NULL;
> }
> } else {
> uint32_t desc_sgpr_idx = *user_sgpr_idx;
> set_userdata_location_shader(ctx,
> AC_UD_INDIRECT_DESCRIPTOR_SETS, user_sgpr_idx, 2);
>
> for (unsigned i = 0; i < num_sets; ++i) {
> if
> (ctx->options->layout->set[i].layout->shader_stages & stage_mask) {
>
> set_userdata_location_indirect(&ctx->shader_info->user_sgprs_locs.descriptor_sets[i],
> desc_sgpr_idx, 2, i * 8);
> - ctx->descriptor_sets[i] =
> ac_build_load_to_sgpr(&ctx->ac, desc_sets, LLVMConstInt(ctx->i32, i, false));
> + ctx->descriptor_sets[i] =
> ac_build_load_to_sgpr(&ctx->ac, desc_sets, LLVMConstInt(ctx->ac.i32, i,
> false));
>
> } else
> ctx->descriptor_sets[i] = NULL;
> }
> ctx->shader_info->need_indirect_descriptor_sets = true;
> }
>
> if (ctx->shader_info->info.needs_push_constants) {
> set_userdata_location_shader(ctx, AC_UD_PUSH_CONSTANTS,
> user_sgpr_idx, 2);
> }
> @@ -686,24 +685,24 @@ radv_define_common_user_sgprs_phase2(struct
> nir_to_llvm_context *ctx,
> static void
> radv_define_vs_user_sgprs_phase1(struct nir_to_llvm_context *ctx,
> gl_shader_stage stage,
> bool has_previous_stage,
> gl_shader_stage previous_stage,
> struct arg_info *args)
> {
> if (!ctx->is_gs_copy_shader && (stage == MESA_SHADER_VERTEX ||
> (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
> if (ctx->shader_info->info.vs.has_vertex_buffers)
> add_user_sgpr_argument(args, const_array(ctx->v4i32,
> 16), &ctx->vertex_buffers); /* vertex buffers */
> - add_user_sgpr_argument(args, ctx->i32,
> &ctx->abi.base_vertex); // base vertex
> - add_user_sgpr_argument(args, ctx->i32,
> &ctx->abi.start_instance);// start instance
> + add_user_sgpr_argument(args, ctx->ac.i32,
> &ctx->abi.base_vertex); // base vertex
> + add_user_sgpr_argument(args, ctx->ac.i32,
> &ctx->abi.start_instance);// start instance
> if (ctx->shader_info->info.vs.needs_draw_id)
> - add_user_sgpr_argument(args, ctx->i32,
> &ctx->abi.draw_id); // draw id
> + add_user_sgpr_argument(args, ctx->ac.i32,
> &ctx->abi.draw_id); // draw id
> }
> }
>
> static void
> radv_define_vs_user_sgprs_phase2(struct nir_to_llvm_context *ctx,
> gl_shader_stage stage,
> bool has_previous_stage,
> gl_shader_stage previous_stage,
> uint8_t *user_sgpr_idx)
> {
> @@ -733,179 +732,179 @@ static void create_function(struct
> nir_to_llvm_context *ctx,
> allocate_user_sgprs(ctx, &user_sgpr_info);
>
> if (user_sgpr_info.need_ring_offsets &&
> !ctx->options->supports_spill) {
> add_user_sgpr_argument(&args, const_array(ctx->v4i32, 16),
> &ctx->ring_offsets); /* address of rings */
> }
>
> switch (stage) {
> case MESA_SHADER_COMPUTE:
> radv_define_common_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
> 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);
> + add_user_sgpr_argument(&args,
> LLVMVectorType(ctx->ac.i32, ctx->shader_info->info.cs.grid_components_used),
> &ctx->num_work_groups); /* grid size */
> + add_sgpr_argument(&args, LLVMVectorType(ctx->ac.i32, 3),
> &ctx->workgroup_ids);
> + add_sgpr_argument(&args, ctx->ac.i32, &ctx->tg_size);
> + add_vgpr_argument(&args, LLVMVectorType(ctx->ac.i32, 3),
> &ctx->local_invocation_ids);
> break;
> case MESA_SHADER_VERTEX:
> radv_define_common_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
> radv_define_vs_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &args);
> if (ctx->shader_info->info.needs_multiview_view_index ||
> (!ctx->options->key.vs.as_es && !ctx->options->key.vs.as_ls &&
> ctx->options->key.has_multiview_view_index))
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->view_index);
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->view_index);
> if (ctx->options->key.vs.as_es)
> - add_sgpr_argument(&args, ctx->i32,
> &ctx->es2gs_offset); // es2gs offset
> + add_sgpr_argument(&args, ctx->ac.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->abi.vertex_id); //
> vertex id
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->ls_out_layout); // ls out layout
> + add_vgpr_argument(&args, ctx->ac.i32, &ctx->abi.vertex_id);
> // 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->abi.instance_id); // instance id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->rel_auto_id); // rel auto id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->vs_prim_id); // vs prim id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->abi.instance_id); // instance id
> }
> break;
> case MESA_SHADER_TESS_CTRL:
> if (has_previous_stage) {
> // First 6 system regs
> - add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); //
> param oc lds
> - add_sgpr_argument(&args, ctx->i32,
> &ctx->merged_wave_info); // merged wave info
> - add_sgpr_argument(&args, ctx->i32,
> &ctx->tess_factor_offset); // tess factor offset
> + add_sgpr_argument(&args, ctx->ac.i32, &ctx->oc_lds);
> // param oc lds
> + add_sgpr_argument(&args, ctx->ac.i32,
> &ctx->merged_wave_info); // merged wave info
> + add_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tess_factor_offset); // tess factor offset
>
> - add_sgpr_argument(&args, ctx->i32, NULL); // scratch
> offset
> - add_sgpr_argument(&args, ctx->i32, NULL); // unknown
> - add_sgpr_argument(&args, ctx->i32, NULL); // unknown
> + add_sgpr_argument(&args, ctx->ac.i32, NULL); //
> scratch offset
> + add_sgpr_argument(&args, ctx->ac.i32, NULL); //
> unknown
> + add_sgpr_argument(&args, ctx->ac.i32, NULL); //
> unknown
>
> radv_define_common_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
> radv_define_vs_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &args);
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->ls_out_layout); // ls out layout
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->ls_out_layout); // ls out layout
>
> - 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_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_offchip_layout); // tcs offchip layout
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_out_offsets); // tcs out offsets
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_out_layout); // tcs out layout
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_in_layout); // tcs in layout
> if (ctx->shader_info->info.needs_multiview_view_index)
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->view_index);
> -
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->tcs_patch_id); // patch id
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->tcs_rel_ids); // rel ids;
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->abi.vertex_id); // vertex 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->abi.instance_id); // instance id
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->view_index);
> +
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_patch_id); // patch id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_rel_ids); // rel ids;
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->abi.vertex_id); // vertex id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->rel_auto_id); // rel auto id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->vs_prim_id); // vs prim id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->abi.instance_id); // instance id
> } else {
> radv_define_common_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
> - 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_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_offchip_layout); // tcs offchip layout
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_out_offsets); // tcs out offsets
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_out_layout); // tcs out layout
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_in_layout); // tcs in layout
> if (ctx->shader_info->info.needs_multiview_view_index)
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->view_index);
> - 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;
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->view_index);
> + add_sgpr_argument(&args, ctx->ac.i32, &ctx->oc_lds);
> // param oc lds
> + add_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tess_factor_offset); // tess factor offset
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_patch_id); // patch id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_rel_ids); // rel ids;
> }
> break;
> case MESA_SHADER_TESS_EVAL:
> radv_define_common_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->tcs_offchip_layout); // tcs offchip layout
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_offchip_layout); // tcs offchip layout
> if (ctx->shader_info->info.needs_multiview_view_index ||
> (!ctx->options->key.tes.as_es && ctx->options->key.has_multiview_view_index))
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->view_index);
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->view_index);
> if (ctx->options->key.tes.as_es) {
> - 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
> + add_sgpr_argument(&args, ctx->ac.i32, &ctx->oc_lds);
> // OC LDS
> + add_sgpr_argument(&args, ctx->ac.i32, NULL); //
> + add_sgpr_argument(&args, ctx->ac.i32,
> &ctx->es2gs_offset); // es2gs offset
> } else {
> - add_sgpr_argument(&args, ctx->i32, NULL); //
> - add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); //
> OC LDS
> + add_sgpr_argument(&args, ctx->ac.i32, NULL); //
> + add_sgpr_argument(&args, ctx->ac.i32, &ctx->oc_lds);
> // OC LDS
> }
> 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
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->tes_rel_patch_id); // tes rel patch id
> + add_vgpr_argument(&args, ctx->ac.i32, &ctx->tes_patch_id); //
> tes patch id
> break;
> case MESA_SHADER_GEOMETRY:
> if (has_previous_stage) {
> // First 6 system regs
> - add_sgpr_argument(&args, ctx->i32,
> &ctx->gs2vs_offset); // tess factor offset
> - add_sgpr_argument(&args, ctx->i32,
> &ctx->merged_wave_info); // merged wave info
> - add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); //
> param oc lds
> + add_sgpr_argument(&args, ctx->ac.i32,
> &ctx->gs2vs_offset); // tess factor offset
> + add_sgpr_argument(&args, ctx->ac.i32,
> &ctx->merged_wave_info); // merged wave info
> + add_sgpr_argument(&args, ctx->ac.i32, &ctx->oc_lds);
> // param oc lds
>
> - add_sgpr_argument(&args, ctx->i32, NULL); // scratch
> offset
> - add_sgpr_argument(&args, ctx->i32, NULL); // unknown
> - add_sgpr_argument(&args, ctx->i32, NULL); // unknown
> + add_sgpr_argument(&args, ctx->ac.i32, NULL); //
> scratch offset
> + add_sgpr_argument(&args, ctx->ac.i32, NULL); //
> unknown
> + add_sgpr_argument(&args, ctx->ac.i32, NULL); //
> unknown
>
> radv_define_common_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
> if (previous_stage == MESA_SHADER_TESS_EVAL)
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->tcs_offchip_layout); // tcs offchip layout
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->tcs_offchip_layout); // tcs offchip layout
> else
> radv_define_vs_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &args);
> - 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_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->gsvs_ring_stride); // gsvs stride
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->gsvs_num_entries); // gsvs num entires
> if (ctx->shader_info->info.needs_multiview_view_index)
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->view_index);
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->view_index);
>
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->gs_vtx_offset[0]); // vtx01
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->gs_vtx_offset[2]); // vtx23
> - add_vgpr_argument(&args, ctx->i32, &ctx->gs_prim_id);
> // prim id
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->gs_invocation_id);
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->gs_vtx_offset[4]);
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[0]); // vtx01
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[2]); // vtx23
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_prim_id); // prim id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_invocation_id);
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[4]);
>
> if (previous_stage == MESA_SHADER_VERTEX) {
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->abi.vertex_id); // vertex 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->abi.instance_id); // instance id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->abi.vertex_id); // vertex id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->rel_auto_id); // rel auto id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->vs_prim_id); // vs prim id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->abi.instance_id); // instance id
> } else {
> 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
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->tes_rel_patch_id); // tes rel patch id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->tes_patch_id); // tes patch id
> }
> } else {
> radv_define_common_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
> radv_define_vs_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &args);
> - 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_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->gsvs_ring_stride); // gsvs stride
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->gsvs_num_entries); // gsvs num entires
> if (ctx->shader_info->info.needs_multiview_view_index)
> - add_user_sgpr_argument(&args, ctx->i32,
> &ctx->view_index);
> - 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);
> + add_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->view_index);
> + add_sgpr_argument(&args, ctx->ac.i32,
> &ctx->gs2vs_offset); // gs2vs offset
> + add_sgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_wave_id); // wave id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[0]); // vtx0
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[1]); // vtx1
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_prim_id); // prim id
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[2]);
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[3]);
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[4]);
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_vtx_offset[5]);
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->gs_invocation_id);
> }
> break;
> case MESA_SHADER_FRAGMENT:
> radv_define_common_user_sgprs_phase1(ctx, stage,
> has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
> if (ctx->shader_info->info.ps.needs_sample_positions)
> - 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_user_sgpr_argument(&args, ctx->ac.i32,
> &ctx->sample_pos_offset); /* sample position offset */
> + add_sgpr_argument(&args, ctx->ac.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->abi.frag_pos[0]);
> /* pos x float */
> add_vgpr_argument(&args, ctx->f32, &ctx->abi.frag_pos[1]);
> /* pos y float */
> add_vgpr_argument(&args, ctx->f32, &ctx->abi.frag_pos[2]);
> /* pos z float */
> add_vgpr_argument(&args, ctx->f32, &ctx->abi.frag_pos[3]);
> /* pos w float */
> - add_vgpr_argument(&args, ctx->i32, &ctx->abi.front_face); /*
> front face */
> - add_vgpr_argument(&args, ctx->i32, &ctx->abi.ancillary); /*
> ancillary */
> - add_vgpr_argument(&args, ctx->i32,
> &ctx->abi.sample_coverage); /* sample coverage */
> - add_vgpr_argument(&args, ctx->i32, NULL); /* fixed pt */
> + add_vgpr_argument(&args, ctx->ac.i32, &ctx->abi.front_face);
> /* front face */
> + add_vgpr_argument(&args, ctx->ac.i32, &ctx->abi.ancillary);
> /* ancillary */
> + add_vgpr_argument(&args, ctx->ac.i32,
> &ctx->abi.sample_coverage); /* sample coverage */
> + add_vgpr_argument(&args, ctx->ac.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, &args,
> ctx->max_workgroup_size,
> ctx->options->unsafe_math);
> set_llvm_calling_convention(ctx->main_function, stage);
> @@ -995,26 +994,25 @@ static void create_function(struct nir_to_llvm_context
> *ctx,
>
> ctx->shader_info->num_user_sgprs = user_sgpr_idx;
> }
>
> static void setup_types(struct nir_to_llvm_context *ctx)
> {
> ctx->voidt = LLVMVoidTypeInContext(ctx->context);
> ctx->i1 = LLVMIntTypeInContext(ctx->context, 1);
> ctx->i8 = LLVMIntTypeInContext(ctx->context, 8);
> ctx->i16 = LLVMIntTypeInContext(ctx->context, 16);
> - ctx->i32 = LLVMIntTypeInContext(ctx->context, 32);
> ctx->i64 = LLVMIntTypeInContext(ctx->context, 64);
> - ctx->v2i32 = LLVMVectorType(ctx->i32, 2);
> - ctx->v3i32 = LLVMVectorType(ctx->i32, 3);
> - ctx->v4i32 = LLVMVectorType(ctx->i32, 4);
> - ctx->v8i32 = LLVMVectorType(ctx->i32, 8);
> + ctx->v2i32 = LLVMVectorType(ctx->ac.i32, 2);
> + ctx->v3i32 = LLVMVectorType(ctx->ac.i32, 3);
> + ctx->v4i32 = LLVMVectorType(ctx->ac.i32, 4);
> + ctx->v8i32 = LLVMVectorType(ctx->ac.i32, 8);
> ctx->f32 = LLVMFloatTypeInContext(ctx->context);
> ctx->f16 = LLVMHalfTypeInContext(ctx->context);
> ctx->f64 = LLVMDoubleTypeInContext(ctx->context);
> ctx->v2f32 = LLVMVectorType(ctx->f32, 2);
> ctx->v4f32 = LLVMVectorType(ctx->f32, 4);
>
> ctx->uniform_md_kind =
> LLVMGetMDKindIDInContext(ctx->context, "amdgpu.uniform", 14);
> ctx->empty_md = LLVMMDNodeInContext(ctx->context, NULL, 0);
> }
> @@ -1337,39 +1335,39 @@ static LLVMValueRef emit_f2f16(struct
> nir_to_llvm_context *ctx,
> LLVMValueRef result;
> LLVMValueRef cond = NULL;
>
> src0 = ac_to_float(&ctx->ac, src0);
> result = LLVMBuildFPTrunc(ctx->builder, src0, ctx->f16, "");
>
> if (ctx->options->chip_class >= VI) {
> LLVMValueRef args[2];
> /* Check if the result is a denormal - and flush to 0 if so.
> */
> args[0] = result;
> - args[1] = LLVMConstInt(ctx->i32, N_SUBNORMAL | P_SUBNORMAL,
> false);
> + args[1] = LLVMConstInt(ctx->ac.i32, N_SUBNORMAL |
> P_SUBNORMAL, false);
> cond = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.class.f16",
> ctx->i1, args, 2, AC_FUNC_ATTR_READNONE);
> }
>
> /* need to convert back up to f32 */
> result = LLVMBuildFPExt(ctx->builder, result, ctx->f32, "");
>
> if (ctx->options->chip_class >= VI)
> result = LLVMBuildSelect(ctx->builder, cond, ctx->ac.f32_0,
> result, "");
> else {
> /* for SI/CIK */
> /* 0x38800000 is smallest half float value (2^-14) in 32-bit
> float,
> * so compare the result and flush to 0 if it's smaller.
> */
> LLVMValueRef temp, cond2;
> temp = emit_intrin_1f_param(&ctx->ac, "llvm.fabs",
> ctx->f32, result);
> cond = LLVMBuildFCmp(ctx->builder, LLVMRealUGT,
> - LLVMBuildBitCast(ctx->builder,
> LLVMConstInt(ctx->i32, 0x38800000, false), ctx->f32, ""),
> + LLVMBuildBitCast(ctx->builder,
> LLVMConstInt(ctx->ac.i32, 0x38800000, false), ctx->f32, ""),
> temp, "");
> cond2 = LLVMBuildFCmp(ctx->builder, LLVMRealUNE,
> temp, ctx->ac.f32_0, "");
> cond = LLVMBuildAnd(ctx->builder, cond, cond2, "");
> result = LLVMBuildSelect(ctx->builder, cond, ctx->ac.f32_0,
> result, "");
> }
> return result;
> }
>
> static LLVMValueRef emit_umul_high(struct ac_llvm_context *ctx,
> @@ -2236,41 +2234,41 @@ static LLVMValueRef
> visit_vulkan_resource_index(struct nir_to_llvm_context *ctx,
> struct radv_descriptor_set_layout *layout =
> pipeline_layout->set[desc_set].layout;
> unsigned base_offset = layout->binding[binding].offset;
> LLVMValueRef offset, stride;
>
> if (layout->binding[binding].type ==
> VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC ||
> layout->binding[binding].type ==
> VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC) {
> unsigned idx =
> pipeline_layout->set[desc_set].dynamic_offset_start +
> layout->binding[binding].dynamic_offset_offset;
> desc_ptr = ctx->push_constants;
> base_offset = pipeline_layout->push_constant_size + 16 * idx;
> - stride = LLVMConstInt(ctx->i32, 16, false);
> + stride = LLVMConstInt(ctx->ac.i32, 16, false);
> } else
> - stride = LLVMConstInt(ctx->i32,
> layout->binding[binding].size, false);
> + stride = LLVMConstInt(ctx->ac.i32,
> layout->binding[binding].size, false);
>
> - offset = LLVMConstInt(ctx->i32, base_offset, false);
> + offset = LLVMConstInt(ctx->ac.i32, base_offset, false);
> index = LLVMBuildMul(ctx->builder, index, stride, "");
> offset = LLVMBuildAdd(ctx->builder, offset, index, "");
>
> desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset);
> desc_ptr = cast_ptr(ctx, desc_ptr, ctx->v4i32);
> LLVMSetMetadata(desc_ptr, ctx->uniform_md_kind, ctx->empty_md);
>
> return LLVMBuildLoad(ctx->builder, desc_ptr, "");
> }
>
> static LLVMValueRef visit_load_push_constant(struct nir_to_llvm_context *ctx,
> nir_intrinsic_instr *instr)
> {
> LLVMValueRef ptr, addr;
>
> - addr = LLVMConstInt(ctx->i32, nir_intrinsic_base(instr), 0);
> + addr = LLVMConstInt(ctx->ac.i32, nir_intrinsic_base(instr), 0);
> addr = LLVMBuildAdd(ctx->builder, addr, get_src(ctx->nir,
> instr->src[0]), "");
>
> ptr = ac_build_gep0(&ctx->ac, ctx->push_constants, addr);
> ptr = cast_ptr(ctx, ptr, get_def_type(ctx->nir, &instr->dest.ssa));
>
> return LLVMBuildLoad(ctx->builder, ptr, "");
> }
>
> static LLVMValueRef visit_get_buffer_size(struct ac_nir_context *ctx,
> const nir_intrinsic_instr *instr)
> @@ -2620,21 +2618,21 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct
> nir_to_llvm_context *ctx,
> {
> LLVMValueRef base_addr, vertices_per_patch, num_patches,
> total_vertices;
> LLVMValueRef param_stride, constant16;
> LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
>
> vertices_per_patch = unpack_param(&ctx->ac, ctx->tcs_offchip_layout,
> 9, 6);
> num_patches = unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 0, 9);
> total_vertices = LLVMBuildMul(ctx->builder, vertices_per_patch,
> num_patches, "");
>
> - constant16 = LLVMConstInt(ctx->i32, 16, false);
> + constant16 = LLVMConstInt(ctx->ac.i32, 16, false);
> if (vertex_index) {
> base_addr = LLVMBuildMul(ctx->builder, rel_patch_id,
> vertices_per_patch, "");
>
> base_addr = LLVMBuildAdd(ctx->builder, base_addr,
> vertex_index, "");
>
> param_stride = total_vertices;
> } else {
> base_addr = rel_patch_id;
> @@ -2660,26 +2658,26 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct
> nir_to_llvm_context *ctx,
> static LLVMValueRef get_tcs_tes_buffer_address_params(struct
> nir_to_llvm_context *ctx,
> unsigned param,
> unsigned const_index,
> bool is_compact,
> LLVMValueRef
> vertex_index,
> LLVMValueRef
> indir_index)
> {
> LLVMValueRef param_index;
>
> if (indir_index)
> - param_index = LLVMBuildAdd(ctx->builder,
> LLVMConstInt(ctx->i32, param, false),
> + param_index = LLVMBuildAdd(ctx->builder,
> LLVMConstInt(ctx->ac.i32, param, false),
> indir_index, "");
> else {
> if (const_index && !is_compact)
> param += const_index;
> - param_index = LLVMConstInt(ctx->i32, param, false);
> + param_index = LLVMConstInt(ctx->ac.i32, param, false);
> }
> return get_tcs_tes_buffer_address(ctx, vertex_index, param_index);
> }
>
> static void
> mark_tess_output(struct nir_to_llvm_context *ctx,
> bool is_patch, uint32_t param)
>
> {
> if (is_patch) {
> @@ -2703,31 +2701,31 @@ get_dw_address(struct nir_to_llvm_context *ctx,
> if (vertex_index) {
> dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
> LLVMBuildMul(ctx->builder,
> vertex_index,
> stride, ""), "");
> }
>
> if (indir_index)
> dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
> LLVMBuildMul(ctx->builder, indir_index,
> - LLVMConstInt(ctx->i32, 4,
> false), ""), "");
> + LLVMConstInt(ctx->ac.i32,
> 4, false), ""), "");
> else if (const_index && !compact_const_index)
> dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
> - LLVMConstInt(ctx->i32, const_index,
> false), "");
> + LLVMConstInt(ctx->ac.i32, const_index,
> false), "");
>
> dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
> - LLVMConstInt(ctx->i32, param * 4, false), "");
> + LLVMConstInt(ctx->ac.i32, param * 4, false),
> "");
>
> if (const_index && compact_const_index)
> dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
> - LLVMConstInt(ctx->i32, const_index,
> false), "");
> + LLVMConstInt(ctx->ac.i32, const_index,
> false), "");
> return dw_addr;
> }
>
> static LLVMValueRef
> build_varying_gather_values(struct ac_llvm_context *ctx, LLVMValueRef
> *values,
> unsigned value_count, unsigned component)
> {
> LLVMValueRef vec = NULL;
>
> if (value_count == 1) {
> @@ -2909,21 +2907,21 @@ load_tes_input(struct nir_to_llvm_context *ctx,
> if (instr->variables[0]->var->data.location ==
> VARYING_SLOT_CLIP_DIST0 &&
> is_compact && const_index > 3) {
> const_index -= 3;
> param++;
> }
>
> unsigned comp = instr->variables[0]->var->data.location_frac;
> buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index,
> is_compact,
> vertex_index, indir_index);
>
> - LLVMValueRef comp_offset = LLVMConstInt(ctx->i32, comp * 4, false);
> + LLVMValueRef comp_offset = LLVMConstInt(ctx->ac.i32, comp * 4, false);
> buf_addr = LLVMBuildAdd(ctx->builder, buf_addr, comp_offset, "");
>
> result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip,
> instr->num_components, NULL,
> buf_addr, ctx->oc_lds, is_compact ? (4
> * const_index) : 0, 1, 0, true, false);
> result = trim_vector(&ctx->ac, result, instr->num_components);
> result = LLVMBuildBitCast(ctx->builder, result,
> get_def_type(ctx->nir, &instr->dest.ssa), "");
> return result;
> }
>
> static LLVMValueRef
> @@ -2935,44 +2933,44 @@ load_gs_input(struct nir_to_llvm_context *ctx,
> LLVMValueRef args[9];
> unsigned param, vtx_offset_param;
> LLVMValueRef value[4], result;
> unsigned vertex_index;
> get_deref_offset(ctx->nir, instr->variables[0],
> false, &vertex_index, NULL,
> &const_index, &indir_index);
> vtx_offset_param = vertex_index;
> assert(vtx_offset_param < 6);
> vtx_offset = LLVMBuildMul(ctx->builder,
> ctx->gs_vtx_offset[vtx_offset_param],
> - LLVMConstInt(ctx->i32, 4, false), "");
> + LLVMConstInt(ctx->ac.i32, 4, false), "");
>
> param =
> shader_io_get_unique_index(instr->variables[0]->var->data.location);
>
> unsigned comp = instr->variables[0]->var->data.location_frac;
> for (unsigned i = comp; i < instr->num_components + comp; i++) {
> if (ctx->ac.chip_class >= GFX9) {
> LLVMValueRef dw_addr =
> ctx->gs_vtx_offset[vtx_offset_param];
> dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
> LLVMConstInt(ctx->ac.i32,
> param * 4 + i + const_index, 0), "");
> value[i] = ac_lds_load(&ctx->ac, dw_addr);
> } else {
> args[0] = ctx->esgs_ring;
> args[1] = vtx_offset;
> - args[2] = LLVMConstInt(ctx->i32, (param * 4 + i +
> const_index) * 256, false);
> + args[2] = LLVMConstInt(ctx->ac.i32, (param * 4 + i +
> const_index) * 256, false);
> args[3] = ctx->ac.i32_0;
> args[4] = ctx->ac.i32_1; /* OFFEN */
> args[5] = ctx->ac.i32_0; /* IDXEN */
> args[6] = ctx->ac.i32_1; /* GLC */
> args[7] = ctx->ac.i32_0; /* SLC */
> args[8] = ctx->ac.i32_0; /* TFE */
>
> value[i] = ac_build_intrinsic(&ctx->ac,
> "llvm.SI.buffer.load.dword.i32.i32",
> - ctx->i32, args, 9,
> + ctx->ac.i32, args, 9,
> AC_FUNC_ATTR_READONLY |
> AC_FUNC_ATTR_LEGACY);
> }
> }
> result = build_varying_gather_values(&ctx->ac, value,
> instr->num_components, comp);
>
> return result;
> }
>
> static LLVMValueRef
> @@ -3685,21 +3683,21 @@ static LLVMValueRef visit_image_size(struct
> ac_nir_context *ctx,
> }
>
> #define NOOP_WAITCNT 0xf7f
> #define LGKM_CNT 0x07f
> #define VM_CNT 0xf70
>
> static void emit_waitcnt(struct nir_to_llvm_context *ctx,
> unsigned simm16)
> {
> LLVMValueRef args[1] = {
> - LLVMConstInt(ctx->i32, simm16, false),
> + LLVMConstInt(ctx->ac.i32, simm16, false),
> };
> ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.s.waitcnt",
> ctx->voidt, args, 1, 0);
> }
>
> static void emit_barrier(struct nir_to_llvm_context *ctx)
> {
> /* SI only (thanks to a hw bug workaround):
> * The real barrier instruction isn’t needed, because an entire patch
> * always fits into a single wave.
> @@ -3723,21 +3721,21 @@ static void emit_discard_if(struct ac_nir_context
> *ctx,
> ctx->ac.i32_0, "");
> ac_build_kill_if_false(&ctx->ac, cond);
> }
>
> static LLVMValueRef
> visit_load_local_invocation_index(struct nir_to_llvm_context *ctx)
> {
> LLVMValueRef result;
> LLVMValueRef thread_id = ac_get_thread_id(&ctx->ac);
> result = LLVMBuildAnd(ctx->builder, ctx->tg_size,
> - LLVMConstInt(ctx->i32, 0xfc0, false), "");
> + LLVMConstInt(ctx->ac.i32, 0xfc0, false), "");
>
> return LLVMBuildAdd(ctx->builder, result, thread_id, "");
> }
>
> static LLVMValueRef visit_var_atomic(struct nir_to_llvm_context *ctx,
> const nir_intrinsic_instr *instr)
> {
> LLVMValueRef ptr, result;
> LLVMValueRef src = get_src(ctx->nir, instr->src[0]);
> ptr = build_gep_for_deref(ctx->nir, instr->variables[0]);
> @@ -3819,21 +3817,21 @@ static LLVMValueRef lookup_interp_param(struct
> nir_to_llvm_context *ctx,
> return ctx->linear_sample;
> break;
> }
> return NULL;
> }
>
> static LLVMValueRef load_sample_position(struct nir_to_llvm_context *ctx,
> LLVMValueRef sample_id)
> {
> LLVMValueRef result;
> - LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets,
> LLVMConstInt(ctx->i32, RING_PS_SAMPLE_POSITIONS, false));
> + LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets,
> LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false));
>
> ptr = LLVMBuildBitCast(ctx->builder, ptr,
> const_array(ctx->v2f32, 64), "");
>
> sample_id = LLVMBuildAdd(ctx->builder, sample_id,
> ctx->sample_pos_offset, "");
> result = ac_build_load_invariant(&ctx->ac, ptr, sample_id);
>
> return result;
> }
>
> @@ -3879,78 +3877,78 @@ static LLVMValueRef visit_interp(struct
> nir_to_llvm_context *ctx,
>
> /* fetch sample ID */
> sample_position = load_sample_position(ctx, src0);
>
> src_c0 = LLVMBuildExtractElement(ctx->builder,
> sample_position, ctx->ac.i32_0, "");
> src_c0 = LLVMBuildFSub(ctx->builder, src_c0, halfval, "");
> src_c1 = LLVMBuildExtractElement(ctx->builder,
> sample_position, ctx->ac.i32_1, "");
> src_c1 = LLVMBuildFSub(ctx->builder, src_c1, halfval, "");
> }
> interp_param = lookup_interp_param(ctx,
> instr->variables[0]->var->data.interpolation, location);
> - attr_number = LLVMConstInt(ctx->i32, input_index, false);
> + attr_number = LLVMConstInt(ctx->ac.i32, input_index, false);
>
> if (location == INTERP_CENTER) {
> LLVMValueRef ij_out[2];
> LLVMValueRef ddxy_out = emit_ddxy_interp(ctx->nir,
> interp_param);
>
> /*
> * take the I then J parameters, and the DDX/Y for it, and
> * calculate the IJ inputs for the interpolator.
> * temp1 = ddx * offset/sample.x + I;
> * interp_param.I = ddy * offset/sample.y + temp1;
> * temp1 = ddx * offset/sample.x + J;
> * interp_param.J = ddy * offset/sample.y + temp1;
> */
> for (unsigned i = 0; i < 2; i++) {
> - LLVMValueRef ix_ll = LLVMConstInt(ctx->i32, i, false);
> - LLVMValueRef iy_ll = LLVMConstInt(ctx->i32, i + 2,
> false);
> + LLVMValueRef ix_ll = LLVMConstInt(ctx->ac.i32, i,
> false);
> + LLVMValueRef iy_ll = LLVMConstInt(ctx->ac.i32, i + 2,
> false);
> LLVMValueRef ddx_el =
> LLVMBuildExtractElement(ctx->builder,
>
> ddxy_out, ix_ll, "");
> LLVMValueRef ddy_el =
> LLVMBuildExtractElement(ctx->builder,
>
> ddxy_out, iy_ll, "");
> LLVMValueRef interp_el =
> LLVMBuildExtractElement(ctx->builder,
>
> interp_param, ix_ll, "");
> LLVMValueRef temp1, temp2;
>
> interp_el = LLVMBuildBitCast(ctx->builder, interp_el,
> ctx->f32, "");
>
> temp1 = LLVMBuildFMul(ctx->builder, ddx_el, src_c0,
> "");
> temp1 = LLVMBuildFAdd(ctx->builder, temp1, interp_el,
> "");
>
> temp2 = LLVMBuildFMul(ctx->builder, ddy_el, src_c1,
> "");
> temp2 = LLVMBuildFAdd(ctx->builder, temp2, temp1, "");
>
> ij_out[i] = LLVMBuildBitCast(ctx->builder,
> - temp2, ctx->i32, "");
> + temp2, ctx->ac.i32, "");
> }
> interp_param = ac_build_gather_values(&ctx->ac, ij_out, 2);
>
> }
>
> for (chan = 0; chan < 4; chan++) {
> - LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, false);
> + LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan,
> false);
>
> if (interp_param) {
> interp_param = LLVMBuildBitCast(ctx->builder,
> interp_param,
> LLVMVectorType(ctx->f32, 2), "");
> LLVMValueRef i = LLVMBuildExtractElement(
> ctx->builder, interp_param, ctx->ac.i32_0,
> "");
> LLVMValueRef j = LLVMBuildExtractElement(
> ctx->builder, interp_param, ctx->ac.i32_1,
> "");
>
> result[chan] = ac_build_fs_interp(&ctx->ac,
> llvm_chan,
> attr_number,
> ctx->prim_mask, i,
> j);
> } else {
> result[chan] = ac_build_fs_interp_mov(&ctx->ac,
> -
> LLVMConstInt(ctx->i32, 2, false),
> +
> LLVMConstInt(ctx->ac.i32, 2, false),
> llvm_chan,
> attr_number,
> ctx->prim_mask);
> }
> }
> return build_varying_gather_values(&ctx->ac, result,
> instr->num_components,
>
> instr->variables[0]->var->data.location_frac);
> }
>
> static void
> visit_emit_vertex(struct nir_to_llvm_context *ctx,
> @@ -3965,21 +3963,21 @@ visit_emit_vertex(struct nir_to_llvm_context *ctx,
> gs_next_vertex = LLVMBuildLoad(ctx->builder,
> ctx->gs_next_vertex,
> "");
>
> /* If this thread has already emitted the declared maximum number of
> * vertices, kill it: excessive vertex emissions are not supposed to
> * have any effect, and GS threads have no externally observable
> * effects other than emitting vertices.
> */
> can_emit = LLVMBuildICmp(ctx->builder, LLVMIntULT, gs_next_vertex,
> - LLVMConstInt(ctx->i32,
> ctx->gs_max_out_vertices, false), "");
> + LLVMConstInt(ctx->ac.i32,
> ctx->gs_max_out_vertices, false), "");
> ac_build_kill_if_false(&ctx->ac, can_emit);
>
> /* loop num outputs */
> idx = 0;
> for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
> LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4];
> int length = 4;
> int slot = idx;
> int slot_inc = 1;
>
> @@ -3988,25 +3986,25 @@ visit_emit_vertex(struct nir_to_llvm_context *ctx,
>
> if (i == VARYING_SLOT_CLIP_DIST0) {
> /* pack clip and cull into a single set of slots */
> length = ctx->num_output_clips +
> ctx->num_output_culls;
> if (length > 4)
> slot_inc = 2;
> }
> for (unsigned j = 0; j < length; j++) {
> LLVMValueRef out_val = LLVMBuildLoad(ctx->builder,
> out_ptr[j], "");
> - LLVMValueRef voffset = LLVMConstInt(ctx->i32, (slot *
> 4 + j) * ctx->gs_max_out_vertices, false);
> + LLVMValueRef voffset = LLVMConstInt(ctx->ac.i32,
> (slot * 4 + j) * ctx->gs_max_out_vertices, false);
> voffset = LLVMBuildAdd(ctx->builder, voffset,
> gs_next_vertex, "");
> - voffset = LLVMBuildMul(ctx->builder, voffset,
> LLVMConstInt(ctx->i32, 4, false), "");
> + voffset = LLVMBuildMul(ctx->builder, voffset,
> LLVMConstInt(ctx->ac.i32, 4, false), "");
>
> - out_val = LLVMBuildBitCast(ctx->builder, out_val,
> ctx->i32, "");
> + out_val = LLVMBuildBitCast(ctx->builder, out_val,
> ctx->ac.i32, "");
>
> ac_build_buffer_store_dword(&ctx->ac, ctx->gsvs_ring,
> out_val, 1,
> voffset,
> ctx->gs2vs_offset, 0,
> 1, 1, true, true);
> }
> idx += slot_inc;
> }
>
> gs_next_vertex = LLVMBuildAdd(ctx->builder, gs_next_vertex,
> @@ -4295,36 +4293,36 @@ static LLVMValueRef radv_get_sampler_desc(struct
> ac_shader_abi *abi,
> offset += constant_index * stride;
>
> if (desc_type == AC_DESC_SAMPLER &&
> binding->immutable_samplers_offset &&
> (!index || binding->immutable_samplers_equal)) {
> if (binding->immutable_samplers_equal)
> constant_index = 0;
>
> const uint32_t *samplers = radv_immutable_samplers(layout,
> binding);
>
> LLVMValueRef constants[] = {
> - LLVMConstInt(ctx->i32, samplers[constant_index * 4 +
> 0], 0),
> - LLVMConstInt(ctx->i32, samplers[constant_index * 4 +
> 1], 0),
> - LLVMConstInt(ctx->i32, samplers[constant_index * 4 +
> 2], 0),
> - LLVMConstInt(ctx->i32, samplers[constant_index * 4 +
> 3], 0),
> + LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4
> + 0], 0),
> + LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4
> + 1], 0),
> + LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4
> + 2], 0),
> + LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4
> + 3], 0),
> };
> return ac_build_gather_values(&ctx->ac, constants, 4);
> }
>
> assert(stride % type_size == 0);
>
> if (!index)
> index = ctx->ac.i32_0;
>
> - index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->i32, stride /
> type_size, 0), "");
> + index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->ac.i32, stride
> / type_size, 0), "");
>
> - list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->i32, offset,
> 0));
> + list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->ac.i32,
> offset, 0));
> list = LLVMBuildPointerCast(builder, list, const_array(type, 0), "");
>
> return ac_build_load_to_sgpr(&ctx->ac, list, index);
> }
>
> static LLVMValueRef get_sampler_desc(struct ac_nir_context *ctx,
> const nir_deref_var *deref,
> enum ac_descriptor_type desc_type,
> const nir_tex_instr *tex_instr,
> bool image, bool write)
> @@ -5021,50 +5019,50 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx,
> if (ctx->options->key.vs.instance_rate_inputs & (1u << index)) {
> buffer_index = LLVMBuildAdd(ctx->builder,
> ctx->abi.instance_id,
> ctx->abi.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->abi.vertex_id,
> ctx->abi.base_vertex, "");
>
> for (unsigned i = 0; i < attrib_count; ++i, ++idx) {
> - t_offset = LLVMConstInt(ctx->i32, index + i, false);
> + t_offset = LLVMConstInt(ctx->ac.i32, index + i, false);
>
> t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr,
> t_offset);
>
> input = ac_build_buffer_load_format(&ctx->ac, t_list,
> buffer_index,
> - LLVMConstInt(ctx->i32, 0,
> false),
> + LLVMConstInt(ctx->ac.i32,
> 0, false),
> true);
>
> for (unsigned chan = 0; chan < 4; chan++) {
> - LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan,
> false);
> + LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32,
> chan, false);
> ctx->inputs[radeon_llvm_reg_index_soa(idx, chan)] =
> ac_to_integer(&ctx->ac,
> LLVMBuildExtractElement(ctx->builder,
> input, llvm_chan,
> ""));
> }
> }
> }
>
> static void interp_fs_input(struct nir_to_llvm_context *ctx,
> unsigned attr,
> LLVMValueRef interp_param,
> LLVMValueRef prim_mask,
> LLVMValueRef result[4])
> {
> LLVMValueRef attr_number;
> unsigned chan;
> LLVMValueRef i, j;
> bool interp = interp_param != NULL;
>
> - attr_number = LLVMConstInt(ctx->i32, attr, false);
> + attr_number = LLVMConstInt(ctx->ac.i32, attr, false);
>
> /* fs.constant returns the param from the middle vertex, so it's not
> * really useful for flat shading. It's meant to be used for custom
> * interpolation (but the intrinsic can't fetch from the other two
> * vertices).
> *
> * Luckily, it doesn't matter, because we rely on the FLAT_SHADE state
> * to do the right thing. The only reason we use fs.constant is that
> * fs.interp cannot be used on integers, because they can be equal
> * to NaN.
> @@ -5073,30 +5071,30 @@ static void interp_fs_input(struct
> nir_to_llvm_context *ctx,
> interp_param = LLVMBuildBitCast(ctx->builder, interp_param,
> LLVMVectorType(ctx->f32, 2),
> "");
>
> i = LLVMBuildExtractElement(ctx->builder, interp_param,
> ctx->ac.i32_0, "");
> j = LLVMBuildExtractElement(ctx->builder, interp_param,
> ctx->ac.i32_1, "");
> }
>
> for (chan = 0; chan < 4; chan++) {
> - LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, false);
> + LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan,
> false);
>
> if (interp) {
> result[chan] = ac_build_fs_interp(&ctx->ac,
> llvm_chan,
> attr_number,
> prim_mask, i, j);
> } else {
> result[chan] = ac_build_fs_interp_mov(&ctx->ac,
> -
> LLVMConstInt(ctx->i32, 2, false),
> +
> LLVMConstInt(ctx->ac.i32, 2, false),
> llvm_chan,
> attr_number,
> prim_mask);
> }
> }
> }
>
> static void
> handle_fs_input_decl(struct nir_to_llvm_context *ctx,
> struct nir_variable *variable)
> @@ -5329,21 +5327,21 @@ handle_shader_output_decl(struct ac_nir_context *ctx,
>
> static LLVMTypeRef
> glsl_base_to_llvm_type(struct nir_to_llvm_context *ctx,
> enum glsl_base_type type)
> {
> switch (type) {
> case GLSL_TYPE_INT:
> case GLSL_TYPE_UINT:
> case GLSL_TYPE_BOOL:
> case GLSL_TYPE_SUBROUTINE:
> - return ctx->i32;
> + return ctx->ac.i32;
> case GLSL_TYPE_FLOAT: /* TODO handle mediump */
> return ctx->f32;
> case GLSL_TYPE_INT64:
> case GLSL_TYPE_UINT64:
> return ctx->i64;
> case GLSL_TYPE_DOUBLE:
> return ctx->f64;
> default:
> unreachable("unknown GLSL type");
> }
> @@ -5431,25 +5429,25 @@ emit_float_saturate(struct ac_llvm_context *ctx,
> LLVMValueRef v, float lo, float
> {
> v = ac_to_float(ctx, v);
> v = emit_intrin_2f_param(ctx, "llvm.maxnum", ctx->f32, v,
> LLVMConstReal(ctx->f32, lo));
> return emit_intrin_2f_param(ctx, "llvm.minnum", ctx->f32, v,
> LLVMConstReal(ctx->f32, hi));
> }
>
>
> static LLVMValueRef emit_pack_int16(struct nir_to_llvm_context *ctx,
> LLVMValueRef src0, LLVMValueRef src1)
> {
> - LLVMValueRef const16 = LLVMConstInt(ctx->i32, 16, false);
> + LLVMValueRef const16 = LLVMConstInt(ctx->ac.i32, 16, false);
> LLVMValueRef comp[2];
>
> - comp[0] = LLVMBuildAnd(ctx->builder, src0, LLVMConstInt(ctx-> i32,
> 65535, 0), "");
> - comp[1] = LLVMBuildAnd(ctx->builder, src1, LLVMConstInt(ctx-> i32,
> 65535, 0), "");
> + comp[0] = LLVMBuildAnd(ctx->builder, src0, LLVMConstInt(ctx->ac.i32,
> 65535, 0), "");
> + comp[1] = LLVMBuildAnd(ctx->builder, src1, LLVMConstInt(ctx->ac.i32,
> 65535, 0), "");
> comp[1] = LLVMBuildShl(ctx->builder, comp[1], const16, "");
> return LLVMBuildOr(ctx->builder, comp[0], comp[1], "");
> }
>
> /* Initialize arguments for the shader export intrinsic */
> static void
> si_llvm_init_export_args(struct nir_to_llvm_context *ctx,
> LLVMValueRef *values,
> unsigned target,
> struct ac_export_args *args)
> @@ -5521,72 +5519,72 @@ si_llvm_init_export_args(struct nir_to_llvm_context
> *ctx,
> break;
>
> case V_028714_SPI_SHADER_UNORM16_ABGR:
> for (unsigned chan = 0; chan < 4; chan++) {
> val[chan] = ac_build_clamp(&ctx->ac,
> values[chan]);
> val[chan] = LLVMBuildFMul(ctx->builder,
> val[chan],
>
> LLVMConstReal(ctx->f32, 65535), "");
> val[chan] = LLVMBuildFAdd(ctx->builder,
> val[chan],
>
> LLVMConstReal(ctx->f32, 0.5), "");
> val[chan] = LLVMBuildFPToUI(ctx->builder,
> val[chan],
> - ctx->i32, "");
> + ctx->ac.i32, "");
> }
>
> args->compr = 1;
> args->out[0] = emit_pack_int16(ctx, val[0], val[1]);
> args->out[1] = emit_pack_int16(ctx, val[2], val[3]);
> break;
>
> case V_028714_SPI_SHADER_SNORM16_ABGR:
> for (unsigned chan = 0; chan < 4; chan++) {
> val[chan] = emit_float_saturate(&ctx->ac,
> values[chan], -1, 1);
> val[chan] = LLVMBuildFMul(ctx->builder,
> val[chan],
>
> LLVMConstReal(ctx->f32, 32767), "");
>
> /* If positive, add 0.5, else add -0.5. */
> val[chan] = LLVMBuildFAdd(ctx->builder,
> val[chan],
> LLVMBuildSelect(ctx->builder,
>
> LLVMBuildFCmp(ctx->builder, LLVMRealOGE,
> val[chan],
> ctx->ac.f32_0, ""),
>
> LLVMConstReal(ctx->f32, 0.5),
>
> LLVMConstReal(ctx->f32, -0.5), ""), "");
> - val[chan] = LLVMBuildFPToSI(ctx->builder,
> val[chan], ctx->i32, "");
> + val[chan] = LLVMBuildFPToSI(ctx->builder,
> val[chan], ctx->ac.i32, "");
> }
>
> args->compr = 1;
> args->out[0] = emit_pack_int16(ctx, val[0], val[1]);
> args->out[1] = emit_pack_int16(ctx, val[2], val[3]);
> break;
>
> case V_028714_SPI_SHADER_UINT16_ABGR: {
> - LLVMValueRef max_rgb = LLVMConstInt(ctx->i32,
> + LLVMValueRef max_rgb = LLVMConstInt(ctx->ac.i32,
> is_int8 ? 255 :
> is_int10 ? 1023 : 65535, 0);
> - LLVMValueRef max_alpha = !is_int10 ? max_rgb :
> LLVMConstInt(ctx->i32, 3, 0);
> + LLVMValueRef max_alpha = !is_int10 ? max_rgb :
> LLVMConstInt(ctx->ac.i32, 3, 0);
>
> for (unsigned chan = 0; chan < 4; chan++) {
> val[chan] = ac_to_integer(&ctx->ac,
> values[chan]);
> val[chan] = emit_minmax_int(&ctx->ac,
> LLVMIntULT, val[chan], chan == 3 ? max_alpha : max_rgb);
> }
>
> args->compr = 1;
> args->out[0] = emit_pack_int16(ctx, val[0], val[1]);
> args->out[1] = emit_pack_int16(ctx, val[2], val[3]);
> break;
> }
>
> case V_028714_SPI_SHADER_SINT16_ABGR: {
> - LLVMValueRef max_rgb = LLVMConstInt(ctx->i32,
> + LLVMValueRef max_rgb = LLVMConstInt(ctx->ac.i32,
> is_int8 ? 127 :
> is_int10 ? 511 : 32767, 0);
> - LLVMValueRef min_rgb = LLVMConstInt(ctx->i32,
> + LLVMValueRef min_rgb = LLVMConstInt(ctx->ac.i32,
> is_int8 ? -128 :
> is_int10 ? -512 : -32768, 0);
> LLVMValueRef max_alpha = !is_int10 ? max_rgb :
> ctx->ac.i32_1;
> - LLVMValueRef min_alpha = !is_int10 ? min_rgb :
> LLVMConstInt(ctx->i32, -2, 0);
> + LLVMValueRef min_alpha = !is_int10 ? min_rgb :
> LLVMConstInt(ctx->ac.i32, -2, 0);
>
> /* Clamp. */
> for (unsigned chan = 0; chan < 4; chan++) {
> val[chan] = ac_to_integer(&ctx->ac,
> values[chan]);
> val[chan] = emit_minmax_int(&ctx->ac,
> LLVMIntSLT, val[chan], chan == 3 ? max_alpha : max_rgb);
> val[chan] = emit_minmax_int(&ctx->ac,
> LLVMIntSGT, val[chan], chan == 3 ? min_alpha : min_rgb);
> }
>
> args->compr = 1;
> args->out[0] = emit_pack_int16(ctx, val[0], val[1]);
> @@ -5707,21 +5705,21 @@ handle_vs_outputs_post(struct nir_to_llvm_context
> *ctx,
> if (outinfo->writes_layer == true)
> pos_args[1].out[2] = layer_value;
> if (outinfo->writes_viewport_index == true) {
> if (ctx->options->chip_class >= GFX9) {
> /* GFX9 has the layer in out.z[10:0] and the
> viewport
> * index in out.z[19:16].
> */
> LLVMValueRef v = viewport_index_value;
> v = ac_to_integer(&ctx->ac, v);
> v = LLVMBuildShl(ctx->builder, v,
> - LLVMConstInt(ctx->i32, 16,
> false),
> + LLVMConstInt(ctx->ac.i32,
> 16, false),
> "");
> v = LLVMBuildOr(ctx->builder, v,
> ac_to_integer(&ctx->ac,
> pos_args[1].out[2]), "");
>
> pos_args[1].out[2] = ac_to_float(&ctx->ac, v);
> pos_args[1].enabled_channels |= 1 << 2;
> } else {
> pos_args[1].out[3] = viewport_index_value;
> pos_args[1].enabled_channels |= 1 << 3;
> }
> @@ -5826,47 +5824,47 @@ handle_es_outputs_post(struct nir_to_llvm_context
> *ctx,
> outinfo->esgs_itemsize = (max_output_written + 1) * 16;
>
> if (ctx->ac.chip_class >= GFX9) {
> unsigned itemsize_dw = outinfo->esgs_itemsize / 4;
> LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
> LLVMValueRef wave_idx = ac_build_bfe(&ctx->ac,
> ctx->merged_wave_info,
>
> LLVMConstInt(ctx->ac.i32, 24, false),
>
> LLVMConstInt(ctx->ac.i32, 4, false), false);
> vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx,
> LLVMBuildMul(ctx->ac.builder,
> wave_idx,
> - LLVMConstInt(ctx->i32,
> 64, false), ""), "");
> +
> LLVMConstInt(ctx->ac.i32, 64, false), ""), "");
> lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx,
> - LLVMConstInt(ctx->i32, itemsize_dw,
> 0), "");
> + LLVMConstInt(ctx->ac.i32,
> itemsize_dw, 0), "");
> }
>
> for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
> LLVMValueRef dw_addr;
> LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4];
> int param_index;
> int length = 4;
>
> if (!(ctx->output_mask & (1ull << i)))
> continue;
>
> if (i == VARYING_SLOT_CLIP_DIST0)
> length = ctx->num_output_clips +
> ctx->num_output_culls;
>
> param_index = shader_io_get_unique_index(i);
>
> if (lds_base) {
> dw_addr = LLVMBuildAdd(ctx->builder, lds_base,
> - LLVMConstInt(ctx->i32,
> param_index * 4, false),
> + LLVMConstInt(ctx->ac.i32,
> param_index * 4, false),
> "");
> }
> for (j = 0; j < length; j++) {
> LLVMValueRef out_val = LLVMBuildLoad(ctx->builder,
> out_ptr[j], "");
> - out_val = LLVMBuildBitCast(ctx->builder, out_val,
> ctx->i32, "");
> + out_val = LLVMBuildBitCast(ctx->builder, out_val,
> ctx->ac.i32, "");
>
> if (ctx->ac.chip_class >= GFX9) {
> ac_lds_store(&ctx->ac, dw_addr,
> LLVMBuildLoad(ctx->builder,
> out_ptr[j], ""));
> dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
> ctx->ac.i32_1, "");
> } else {
> ac_build_buffer_store_dword(&ctx->ac,
> ctx->esgs_ring,
> out_val, 1,
> NULL,
> ctx->es2gs_offset,
> @@ -5892,21 +5890,21 @@ handle_ls_outputs_post(struct nir_to_llvm_context
> *ctx)
> if (!(ctx->output_mask & (1ull << i)))
> continue;
>
> if (i == VARYING_SLOT_CLIP_DIST0)
> length = ctx->num_output_clips +
> ctx->num_output_culls;
> int param = shader_io_get_unique_index(i);
> mark_tess_output(ctx, false, param);
> if (length > 4)
> mark_tess_output(ctx, false, param + 1);
> LLVMValueRef dw_addr = LLVMBuildAdd(ctx->builder,
> base_dw_addr,
> - LLVMConstInt(ctx->i32,
> param * 4, false),
> + LLVMConstInt(ctx->ac.i32,
> param * 4, false),
> "");
> for (unsigned j = 0; j < length; j++) {
> ac_lds_store(&ctx->ac, dw_addr,
> LLVMBuildLoad(ctx->builder, out_ptr[j],
> ""));
> dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
> ctx->ac.i32_1, "");
> }
> }
> }
>
> struct ac_build_if_state
> @@ -6037,72 +6035,72 @@ write_tess_factors(struct nir_to_llvm_context *ctx)
> LLVMBuildICmp(ctx->builder, LLVMIntEQ,
> invocation_id, ctx->ac.i32_0, ""));
>
> tess_inner_index =
> shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
> tess_outer_index =
> shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
>
> mark_tess_output(ctx, true, tess_inner_index);
> mark_tess_output(ctx, true, tess_outer_index);
> lds_base = get_tcs_out_current_patch_data_offset(ctx);
> lds_inner = LLVMBuildAdd(ctx->builder, lds_base,
> - LLVMConstInt(ctx->i32, tess_inner_index * 4,
> false), "");
> + LLVMConstInt(ctx->ac.i32, tess_inner_index *
> 4, false), "");
> lds_outer = LLVMBuildAdd(ctx->builder, lds_base,
> - LLVMConstInt(ctx->i32, tess_outer_index * 4,
> false), "");
> + LLVMConstInt(ctx->ac.i32, tess_outer_index *
> 4, false), "");
>
> for (i = 0; i < 4; i++) {
> - inner[i] = LLVMGetUndef(ctx->i32);
> - outer[i] = LLVMGetUndef(ctx->i32);
> + inner[i] = LLVMGetUndef(ctx->ac.i32);
> + outer[i] = LLVMGetUndef(ctx->ac.i32);
> }
>
> // LINES reverseal
> if (ctx->options->key.tcs.primitive_mode == GL_ISOLINES) {
> outer[0] = out[1] = ac_lds_load(&ctx->ac, lds_outer);
> lds_outer = LLVMBuildAdd(ctx->builder, lds_outer,
> - LLVMConstInt(ctx->i32, 1, false),
> "");
> + LLVMConstInt(ctx->ac.i32, 1, false),
> "");
> outer[1] = out[0] = ac_lds_load(&ctx->ac, lds_outer);
> } else {
> for (i = 0; i < outer_comps; i++) {
> outer[i] = out[i] =
> ac_lds_load(&ctx->ac, lds_outer);
> lds_outer = LLVMBuildAdd(ctx->builder, lds_outer,
> - LLVMConstInt(ctx->i32, 1,
> false), "");
> + LLVMConstInt(ctx->ac.i32, 1,
> false), "");
> }
> for (i = 0; i < inner_comps; i++) {
> inner[i] = out[outer_comps+i] =
> ac_lds_load(&ctx->ac, lds_inner);
> lds_inner = LLVMBuildAdd(ctx->builder, lds_inner,
> - LLVMConstInt(ctx->i32, 1,
> false), "");
> + LLVMConstInt(ctx->ac.i32, 1,
> false), "");
> }
> }
>
> /* Convert the outputs to vectors for stores. */
> vec0 = ac_build_gather_values(&ctx->ac, out, MIN2(stride, 4));
> vec1 = NULL;
>
> if (stride > 4)
> vec1 = ac_build_gather_values(&ctx->ac, out + 4, stride - 4);
>
>
> buffer = ctx->hs_ring_tess_factor;
> tf_base = ctx->tess_factor_offset;
> byteoffset = LLVMBuildMul(ctx->builder, rel_patch_id,
> - LLVMConstInt(ctx->i32, 4 * stride, false),
> "");
> + LLVMConstInt(ctx->ac.i32, 4 * stride,
> false), "");
> unsigned tf_offset = 0;
>
> if (ctx->options->chip_class <= VI) {
> ac_nir_build_if(&inner_if_ctx, ctx,
> LLVMBuildICmp(ctx->builder, LLVMIntEQ,
> rel_patch_id, ctx->ac.i32_0,
> ""));
>
> /* Store the dynamic HS control word. */
> ac_build_buffer_store_dword(&ctx->ac, buffer,
> - LLVMConstInt(ctx->i32,
> 0x80000000, false),
> + LLVMConstInt(ctx->ac.i32,
> 0x80000000, false),
> 1, ctx->ac.i32_0, tf_base,
> 0, 1, 0, true, false);
> tf_offset += 4;
>
> ac_nir_build_endif(&inner_if_ctx);
> }
>
> /* Store the tessellation factors. */
> ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
> MIN2(stride, 4), byteoffset, tf_base,
> @@ -6113,32 +6111,32 @@ write_tess_factors(struct nir_to_llvm_context *ctx)
> 16 + tf_offset, 1, 0, true,
> false);
>
> //store to offchip for TES to read - only if TES reads them
> if (ctx->options->key.tcs.tes_reads_tess_factors) {
> LLVMValueRef inner_vec, outer_vec, tf_outer_offset;
> LLVMValueRef tf_inner_offset;
> unsigned param_outer, param_inner;
>
> param_outer =
> shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
> tf_outer_offset = get_tcs_tes_buffer_address(ctx, NULL,
> -
> LLVMConstInt(ctx->i32, param_outer, 0));
> +
> LLVMConstInt(ctx->ac.i32, param_outer, 0));
>
> outer_vec = ac_build_gather_values(&ctx->ac, outer,
>
> util_next_power_of_two(outer_comps));
>
> ac_build_buffer_store_dword(&ctx->ac,
> ctx->hs_ring_tess_offchip, outer_vec,
> outer_comps, tf_outer_offset,
> ctx->oc_lds, 0, 1, 0, true,
> false);
> if (inner_comps) {
> param_inner =
> shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
> tf_inner_offset = get_tcs_tes_buffer_address(ctx,
> NULL,
> -
> LLVMConstInt(ctx->i32, param_inner, 0));
> +
> LLVMConstInt(ctx->ac.i32, param_inner, 0));
>
> inner_vec = inner_comps == 1 ? inner[0] :
> ac_build_gather_values(&ctx->ac, inner,
> inner_comps);
> ac_build_buffer_store_dword(&ctx->ac,
> ctx->hs_ring_tess_offchip, inner_vec,
> inner_comps,
> tf_inner_offset,
> ctx->oc_lds, 0, 1, 0,
> true, false);
> }
> }
> ac_nir_build_endif(&if_ctx);
> }
> @@ -6362,43 +6360,43 @@ ac_nir_eliminate_const_vs_outputs(struct
> nir_to_llvm_context *ctx)
> outinfo->vs_output_param_offset,
> VARYING_SLOT_MAX,
> &outinfo->param_exports);
> }
>
> static void
> ac_setup_rings(struct nir_to_llvm_context *ctx)
> {
> if ((ctx->stage == MESA_SHADER_VERTEX && ctx->options->key.vs.as_es)
> ||
> (ctx->stage == MESA_SHADER_TESS_EVAL &&
> ctx->options->key.tes.as_es)) {
> - ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_ESGS_VS, false));
> + ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_ESGS_VS, false));
> }
>
> if (ctx->is_gs_copy_shader) {
> - ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_GSVS_VS, false));
> + ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_GSVS_VS, false));
> }
> if (ctx->stage == MESA_SHADER_GEOMETRY) {
> LLVMValueRef tmp;
> - ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_ESGS_GS, false));
> - ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_GSVS_GS, false));
> + ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_ESGS_GS, false));
> + ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_GSVS_GS, false));
>
> ctx->gsvs_ring = LLVMBuildBitCast(ctx->builder,
> ctx->gsvs_ring, ctx->v4i32, "");
>
> - ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder,
> ctx->gsvs_ring, ctx->gsvs_num_entries, LLVMConstInt(ctx->i32, 2, false), "");
> + ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder,
> ctx->gsvs_ring, ctx->gsvs_num_entries, LLVMConstInt(ctx->ac.i32, 2, false),
> "");
> tmp = LLVMBuildExtractElement(ctx->builder, ctx->gsvs_ring,
> ctx->ac.i32_1, "");
> tmp = LLVMBuildOr(ctx->builder, tmp, ctx->gsvs_ring_stride,
> "");
> ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder,
> ctx->gsvs_ring, tmp, ctx->ac.i32_1, "");
> }
>
> if (ctx->stage == MESA_SHADER_TESS_CTRL ||
> ctx->stage == MESA_SHADER_TESS_EVAL) {
> - ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_HS_TESS_OFFCHIP, false));
> - ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_HS_TESS_FACTOR, false));
> + ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_OFFCHIP, false));
> + ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(&ctx->ac,
> ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_FACTOR, false));
> }
> }
>
> static unsigned
> ac_nir_get_max_workgroup_size(enum chip_class chip_class,
> const struct nir_shader *nir)
> {
> switch (nir->info.stage) {
> case MESA_SHADER_TESS_CTRL:
> return chip_class >= CIK ? 128 : 64;
> @@ -6554,21 +6552,21 @@ LLVMModuleRef
> ac_translate_nir_to_llvm(LLVMTargetMachineRef tm,
> ac_nir_fixup_ls_hs_input_vgprs(&ctx);
>
> for(int i = 0; i < shader_count; ++i) {
> ctx.stage = shaders[i]->info.stage;
> ctx.output_mask = 0;
> ctx.tess_outputs_written = 0;
> ctx.num_output_clips =
> shaders[i]->info.clip_distance_array_size;
> ctx.num_output_culls =
> shaders[i]->info.cull_distance_array_size;
>
> if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
> - ctx.gs_next_vertex = ac_build_alloca(&ctx.ac,
> ctx.i32, "gs_next_vertex");
> + ctx.gs_next_vertex = ac_build_alloca(&ctx.ac,
> ctx.ac.i32, "gs_next_vertex");
>
> ctx.gs_max_out_vertices =
> shaders[i]->info.gs.vertices_out;
> } else if (shaders[i]->info.stage == MESA_SHADER_TESS_EVAL) {
> ctx.tes_primitive_mode =
> shaders[i]->info.tess.primitive_mode;
> } else if (shaders[i]->info.stage == MESA_SHADER_VERTEX) {
> if (shader_info->info.vs.needs_instance_id) {
> ctx.shader_info->vs.vgpr_comp_cnt =
> MAX2(3,
> ctx.shader_info->vs.vgpr_comp_cnt);
> }
> } else if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT) {
> @@ -6819,21 +6817,21 @@ void ac_compile_nir_shader(LLVMTargetMachineRef tm,
> ac_compile_llvm_module(tm, llvm_module, binary, config, shader_info,
> nir[0]->info.stage, dump_shader, options->supports_spill);
> for (int i = 0; i < nir_count; ++i)
> ac_fill_shader_info(shader_info, nir[i], options);
> }
>
> 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->abi.vertex_id,
> LLVMConstInt(ctx->i32, 4, false), "");
> + args[1] = LLVMBuildMul(ctx->builder, ctx->abi.vertex_id,
> LLVMConstInt(ctx->ac.i32, 4, false), "");
> args[3] = ctx->ac.i32_0;
> args[4] = ctx->ac.i32_1; /* OFFEN */
> args[5] = ctx->ac.i32_0; /* IDXEN */
> args[6] = ctx->ac.i32_1; /* GLC */
> args[7] = ctx->ac.i32_1; /* SLC */
> args[8] = ctx->ac.i32_0; /* TFE */
>
> int idx = 0;
>
> for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
> @@ -6845,27 +6843,27 @@ ac_gs_copy_shader_emit(struct nir_to_llvm_context
> *ctx)
>
> if (i == VARYING_SLOT_CLIP_DIST0) {
> /* unpack clip and cull from a single set of slots */
> length = ctx->num_output_clips +
> ctx->num_output_culls;
> if (length > 4)
> slot_inc = 2;
> }
>
> for (unsigned j = 0; j < length; j++) {
> LLVMValueRef value;
> - args[2] = LLVMConstInt(ctx->i32,
> + args[2] = LLVMConstInt(ctx->ac.i32,
> (slot * 4 + j) *
> ctx->gs_max_out_vertices * 16
> * 4, false);
>
> value = ac_build_intrinsic(&ctx->ac,
>
> "llvm.SI.buffer.load.dword.i32.i32",
> - ctx->i32, args, 9,
> + ctx->ac.i32, args, 9,
> AC_FUNC_ATTR_READONLY |
> AC_FUNC_ATTR_LEGACY);
>
> LLVMBuildStore(ctx->builder,
> ac_to_float(&ctx->ac, value),
> ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)]);
> }
> idx += slot_inc;
> }
> handle_vs_outputs_post(ctx, false, &ctx->shader_info->vs.outinfo);
> }
> --
> 2.14.3
>
> _______________________________________________
> mesa-dev mailing list
> mesa-dev@lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
