- for (int i = 0; i < 3; i++) {
- ctx->abi.workgroup_ids[i] = NULL;
- if (ctx->shader_info->cs.uses_block_id[i]) {
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->abi.workgroup_ids[i]);
- }
- }
-
- if (ctx->shader_info->cs.uses_local_invocation_idx)
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.tg_size);
- add_arg(&args, ARG_VGPR, ctx->ac.v3i32,
- &ctx->abi.local_invocation_ids);
- break;
- case MESA_SHADER_VERTEX:
- declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
- &desc_sets);
-
- declare_vs_specific_input_sgprs(ctx, stage, has_previous_stage,
- previous_stage, &args);
-
- if (needs_view_index)
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->abi.view_index);
- if (ctx->options->key.vs_common_out.as_es) {
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->es2gs_offset);
- } else if (ctx->options->key.vs_common_out.as_ls) {
- /* no extra parameters */
- } else {
- declare_streamout_sgprs(ctx, stage, &args);
- }
-
- declare_vs_input_vgprs(ctx, &args);
- break;
- case MESA_SHADER_TESS_CTRL:
- if (has_previous_stage) {
- // First 6 system regs
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->merged_wave_info);
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->tess_factor_offset);
-
- add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
- add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
- add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
-
- declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
- &desc_sets);
-
- declare_vs_specific_input_sgprs(ctx, stage,
- has_previous_stage,
- previous_stage, &args);
-
- if (needs_view_index)
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->abi.view_index);
-
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->abi.tcs_patch_id);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->abi.tcs_rel_ids);
-
- declare_vs_input_vgprs(ctx, &args);
- } else {
- declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
- &desc_sets);
-
- if (needs_view_index)
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->abi.view_index);
-
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->tess_factor_offset);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->abi.tcs_patch_id);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->abi.tcs_rel_ids);
- }
- break;
- case MESA_SHADER_TESS_EVAL:
- declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
- &desc_sets);
-
- if (needs_view_index)
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->abi.view_index);
-
- if (ctx->options->key.vs_common_out.as_es) {
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
- add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->es2gs_offset);
- } else {
- add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
- declare_streamout_sgprs(ctx, stage, &args);
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
- }
- declare_tes_input_vgprs(ctx, &args);
- break;
- case MESA_SHADER_GEOMETRY:
- if (has_previous_stage) {
- // First 6 system regs
- if (ctx->options->key.vs_common_out.as_ngg) {
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->gs_tg_info);
- } else {
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->gs2vs_offset);
- }
-
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->merged_wave_info);
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
-
- add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
- add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
- add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
-
- declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
- &desc_sets);
-
- if (previous_stage != MESA_SHADER_TESS_EVAL) {
- declare_vs_specific_input_sgprs(ctx, stage,
- has_previous_stage,
- previous_stage,
- &args);
- }
-
- if (needs_view_index)
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->abi.view_index);
-
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[0]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[2]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->abi.gs_prim_id);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->abi.gs_invocation_id);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[4]);
-
- if (previous_stage == MESA_SHADER_VERTEX) {
- declare_vs_input_vgprs(ctx, &args);
- } else {
- declare_tes_input_vgprs(ctx, &args);
- }
- } else {
- declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
- &desc_sets);
-
- if (needs_view_index)
- add_arg(&args, ARG_SGPR, ctx->ac.i32,
- &ctx->abi.view_index);
-
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs2vs_offset);
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs_wave_id);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[0]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[1]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->abi.gs_prim_id);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[2]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[3]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[4]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->gs_vtx_offset[5]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32,
- &ctx->abi.gs_invocation_id);
- }
- break;
- case MESA_SHADER_FRAGMENT:
- declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
- &desc_sets);
-
- add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.prim_mask);
- add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_sample);
- add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_center);
- add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_centroid);
- add_arg(&args, ARG_VGPR, ctx->ac.v3i32, NULL); /* persp pull model */
- add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_sample);
- add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_center);
- add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_centroid);
- add_arg(&args, ARG_VGPR, ctx->ac.f32, NULL); /* line stipple tex */
- add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[0]);
- add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[1]);
- add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[2]);
- add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[3]);
- add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.front_face);
- add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.ancillary);
- add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.sample_coverage);
- add_arg(&args, ARG_VGPR, ctx->ac.i32, NULL); /* fixed pt */
- break;
- default:
- unreachable("Shader stage not implemented");
- }
-
- ctx->main_function = create_llvm_function(
- ctx->context, ctx->ac.module, ctx->ac.builder, NULL, 0, &args,
- ctx->max_workgroup_size, ctx->options);
- set_llvm_calling_convention(ctx->main_function, stage);
-
-
- ctx->shader_info->num_input_vgprs = 0;
- ctx->shader_info->num_input_sgprs = ctx->options->supports_spill ? 2 : 0;
-
- ctx->shader_info->num_input_sgprs += args.num_sgprs_used;
-
- if (ctx->stage != MESA_SHADER_FRAGMENT)
- ctx->shader_info->num_input_vgprs = args.num_vgprs_used;
-
- assign_arguments(ctx->main_function, &args);
-
- user_sgpr_idx = 0;
-
- if (ctx->options->supports_spill || user_sgpr_info.need_ring_offsets) {
- set_loc_shader_ptr(ctx, AC_UD_SCRATCH_RING_OFFSETS,
- &user_sgpr_idx);
- if (ctx->options->supports_spill) {
- ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
- LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_CONST),
- NULL, 0, AC_FUNC_ATTR_READNONE);
- ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets,
- ac_array_in_const_addr_space(ctx->ac.v4i32), "");
- }
- }
-
- /* For merged shaders the user SGPRs start at 8, with 8 system SGPRs in front (including
- * the rw_buffers at s0/s1. With user SGPR0 = s8, lets restart the count from 0 */
- if (has_previous_stage)
- user_sgpr_idx = 0;
-
- set_global_input_locs(ctx, &user_sgpr_info, desc_sets, &user_sgpr_idx);
-
- switch (stage) {
- case MESA_SHADER_COMPUTE:
- if (ctx->shader_info->cs.uses_grid_size) {
- set_loc_shader(ctx, AC_UD_CS_GRID_SIZE,
- &user_sgpr_idx, 3);
- }
- break;
- case MESA_SHADER_VERTEX:
- set_vs_specific_input_locs(ctx, stage, has_previous_stage,
- previous_stage, &user_sgpr_idx);
- if (ctx->abi.view_index)
- set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
- break;
- case MESA_SHADER_TESS_CTRL:
- set_vs_specific_input_locs(ctx, stage, has_previous_stage,
- previous_stage, &user_sgpr_idx);
- if (ctx->abi.view_index)
- set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
- break;
- case MESA_SHADER_TESS_EVAL:
- if (ctx->abi.view_index)
- set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
- break;
- case MESA_SHADER_GEOMETRY:
- if (has_previous_stage) {
- if (previous_stage == MESA_SHADER_VERTEX)
- set_vs_specific_input_locs(ctx, stage,
- has_previous_stage,
- previous_stage,
- &user_sgpr_idx);
- }
- if (ctx->abi.view_index)
- set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
- break;
- case MESA_SHADER_FRAGMENT:
- break;
- default:
- unreachable("Shader stage not implemented");
- }
-
- if (stage == MESA_SHADER_TESS_CTRL ||
- (stage == MESA_SHADER_VERTEX && ctx->options->key.vs_common_out.as_ls) ||
- /* GFX9 has the ESGS ring buffer in LDS. */
- (stage == MESA_SHADER_GEOMETRY && has_previous_stage)) {
- ac_declare_lds_as_pointer(&ctx->ac);
- }
-
- ctx->shader_info->num_user_sgprs = user_sgpr_idx;