- set_loc_desc(ctx, i, user_sgpr_idx);
- }
- } else {
- set_loc_shader_ptr(ctx, AC_UD_INDIRECT_DESCRIPTOR_SETS,
- user_sgpr_idx);
-
- while (mask) {
- int i = u_bit_scan(&mask);
-
- ctx->descriptor_sets[i] =
- ac_build_load_to_sgpr(&ctx->ac, desc_sets,
- LLVMConstInt(ctx->ac.i32, i, false));
-
- }
-
- ctx->shader_info->need_indirect_descriptor_sets = true;
- }
-
- if (ctx->shader_info->loads_push_constants) {
- set_loc_shader_ptr(ctx, AC_UD_PUSH_CONSTANTS, user_sgpr_idx);
- }
-
- if (ctx->shader_info->num_inline_push_consts) {
- set_loc_shader(ctx, AC_UD_INLINE_PUSH_CONSTANTS, user_sgpr_idx,
- ctx->shader_info->num_inline_push_consts);
- }
-
- if (ctx->streamout_buffers) {
- set_loc_shader_ptr(ctx, AC_UD_STREAMOUT_BUFFERS,
- user_sgpr_idx);
- }
-}
-
-static void
-set_vs_specific_input_locs(struct radv_shader_context *ctx,
- gl_shader_stage stage, bool has_previous_stage,
- gl_shader_stage previous_stage,
- uint8_t *user_sgpr_idx)
-{
- if (!ctx->is_gs_copy_shader &&
- (stage == MESA_SHADER_VERTEX ||
- (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
- if (ctx->shader_info->vs.has_vertex_buffers) {
- set_loc_shader_ptr(ctx, AC_UD_VS_VERTEX_BUFFERS,
- user_sgpr_idx);
- }
-
- unsigned vs_num = 2;
- if (ctx->shader_info->vs.needs_draw_id)
- vs_num++;
-
- set_loc_shader(ctx, AC_UD_VS_BASE_VERTEX_START_INSTANCE,
- user_sgpr_idx, vs_num);
- }
-}
-
-static void set_llvm_calling_convention(LLVMValueRef func,
- gl_shader_stage stage)
-{
- enum radeon_llvm_calling_convention calling_conv;
-
- switch (stage) {
- case MESA_SHADER_VERTEX:
- case MESA_SHADER_TESS_EVAL:
- calling_conv = RADEON_LLVM_AMDGPU_VS;
- break;
- case MESA_SHADER_GEOMETRY:
- calling_conv = RADEON_LLVM_AMDGPU_GS;
- break;
- case MESA_SHADER_TESS_CTRL:
- calling_conv = RADEON_LLVM_AMDGPU_HS;
- break;
- case MESA_SHADER_FRAGMENT:
- calling_conv = RADEON_LLVM_AMDGPU_PS;
- break;
- case MESA_SHADER_COMPUTE:
- calling_conv = RADEON_LLVM_AMDGPU_CS;
- break;
- default:
- unreachable("Unhandle shader type");
- }
-
- LLVMSetFunctionCallConv(func, calling_conv);
-}
-
-/* Returns whether the stage is a stage that can be directly before the GS */
-static bool is_pre_gs_stage(gl_shader_stage stage)
-{
- return stage == MESA_SHADER_VERTEX || stage == MESA_SHADER_TESS_EVAL;
-}
-
-static void create_function(struct radv_shader_context *ctx,
- gl_shader_stage stage,
- bool has_previous_stage,
- gl_shader_stage previous_stage)
-{
- uint8_t user_sgpr_idx;
- struct user_sgpr_info user_sgpr_info;
- struct arg_info args = {};
- LLVMValueRef desc_sets;
- bool needs_view_index = needs_view_index_sgpr(ctx, stage);
-
- if (ctx->ac.chip_class >= GFX10) {
- if (is_pre_gs_stage(stage) && ctx->options->key.vs_common_out.as_ngg) {
- /* On GFX10, VS is merged into GS for NGG. */
- previous_stage = stage;
- stage = MESA_SHADER_GEOMETRY;
- has_previous_stage = true;
- }
- }
-
- allocate_user_sgprs(ctx, stage, has_previous_stage,
- previous_stage, needs_view_index, &user_sgpr_info);
-
- if (user_sgpr_info.need_ring_offsets && !ctx->options->supports_spill) {
- add_arg(&args, ARG_SGPR, ac_array_in_const_addr_space(ctx->ac.v4i32),
- &ctx->ring_offsets);
- }
-
- switch (stage) {
- case MESA_SHADER_COMPUTE:
- declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
- &desc_sets);
-
- if (ctx->shader_info->cs.uses_grid_size) {
- add_arg(&args, ARG_SGPR, ctx->ac.v3i32,
- &ctx->abi.num_work_groups);
- }
-
- 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);