#include "radv_shader.h"
#include "radv_shader_helper.h"
#include "radv_shader_args.h"
+#include "radv_debug.h"
#include "nir/nir.h"
-#include <llvm-c/Core.h>
-#include <llvm-c/TargetMachine.h>
-#include <llvm-c/Transforms/Scalar.h>
-#include <llvm-c/Transforms/Utils.h>
-
#include "sid.h"
#include "ac_binary.h"
#include "ac_llvm_util.h"
}
}
-static unsigned
-get_tcs_num_patches(struct radv_shader_context *ctx)
-{
- unsigned num_tcs_input_cp = ctx->args->options->key.tcs.input_vertices;
- unsigned num_tcs_output_cp = ctx->shader->info.tess.tcs_vertices_out;
- uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
- uint32_t input_patch_size = ctx->args->options->key.tcs.input_vertices * input_vertex_size;
- uint32_t num_tcs_outputs = util_last_bit64(ctx->args->shader_info->tcs.outputs_written);
- uint32_t num_tcs_patch_outputs = util_last_bit64(ctx->args->shader_info->tcs.patch_outputs_written);
- uint32_t output_vertex_size = num_tcs_outputs * 16;
- uint32_t pervertex_output_patch_size = ctx->shader->info.tess.tcs_vertices_out * output_vertex_size;
- uint32_t output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
- unsigned num_patches;
- unsigned hardware_lds_size;
-
- /* Ensure that we only need one wave per SIMD so we don't need to check
- * resource usage. Also ensures that the number of tcs in and out
- * vertices per threadgroup are at most 256.
- */
- num_patches = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp) * 4;
- /* Make sure that the data fits in LDS. This assumes the shaders only
- * use LDS for the inputs and outputs.
- */
- hardware_lds_size = 32768;
-
- /* Looks like STONEY hangs if we use more than 32 KiB LDS in a single
- * threadgroup, even though there is more than 32 KiB LDS.
- *
- * Test: dEQP-VK.tessellation.shader_input_output.barrier
- */
- if (ctx->args->options->chip_class >= GFX7 && ctx->args->options->family != CHIP_STONEY)
- hardware_lds_size = 65536;
-
- num_patches = MIN2(num_patches, hardware_lds_size / (input_patch_size + output_patch_size));
- /* Make sure the output data fits in the offchip buffer */
- num_patches = MIN2(num_patches, (ctx->args->options->tess_offchip_block_dw_size * 4) / output_patch_size);
- /* Not necessary for correctness, but improves performance. The
- * specific value is taken from the proprietary driver.
- */
- num_patches = MIN2(num_patches, 40);
-
- /* GFX6 bug workaround - limit LS-HS threadgroups to only one wave. */
- if (ctx->args->options->chip_class == GFX6) {
- unsigned one_wave = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp);
- num_patches = MIN2(num_patches, one_wave);
- }
- return num_patches;
-}
-
-static unsigned
-calculate_tess_lds_size(struct radv_shader_context *ctx)
-{
- unsigned num_tcs_input_cp = ctx->args->options->key.tcs.input_vertices;
- unsigned num_tcs_output_cp;
- unsigned num_tcs_outputs, num_tcs_patch_outputs;
- unsigned input_vertex_size, output_vertex_size;
- unsigned input_patch_size, output_patch_size;
- unsigned pervertex_output_patch_size;
- unsigned output_patch0_offset;
- unsigned num_patches;
- unsigned lds_size;
-
- num_tcs_output_cp = ctx->shader->info.tess.tcs_vertices_out;
- num_tcs_outputs = util_last_bit64(ctx->args->shader_info->tcs.outputs_written);
- num_tcs_patch_outputs = util_last_bit64(ctx->args->shader_info->tcs.patch_outputs_written);
-
- input_vertex_size = ctx->tcs_num_inputs * 16;
- output_vertex_size = num_tcs_outputs * 16;
-
- input_patch_size = num_tcs_input_cp * input_vertex_size;
-
- pervertex_output_patch_size = num_tcs_output_cp * output_vertex_size;
- output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
-
- num_patches = ctx->tcs_num_patches;
- output_patch0_offset = input_patch_size * num_patches;
-
- lds_size = output_patch0_offset + output_patch_size * num_patches;
- return lds_size;
-}
-
/* Tessellation shaders pass outputs to the next shader using LDS.
*
* LS outputs = TCS inputs
static LLVMValueRef
create_llvm_function(struct ac_llvm_context *ctx, LLVMModuleRef module,
LLVMBuilderRef builder,
- struct ac_shader_args *args,
+ const struct ac_shader_args *args,
enum ac_llvm_calling_convention convention,
unsigned max_workgroup_size,
const struct radv_nir_compiler_options *options)
return main_function;
}
-
-static void
-set_loc(struct radv_userdata_info *ud_info, uint8_t *sgpr_idx,
- uint8_t num_sgprs)
-{
- ud_info->sgpr_idx = *sgpr_idx;
- ud_info->num_sgprs = num_sgprs;
- *sgpr_idx += num_sgprs;
-}
-
-static void
-set_loc_shader(struct radv_shader_args *args, int idx, uint8_t *sgpr_idx,
- uint8_t num_sgprs)
-{
- struct radv_userdata_info *ud_info =
- &args->shader_info->user_sgprs_locs.shader_data[idx];
- assert(ud_info);
-
- set_loc(ud_info, sgpr_idx, num_sgprs);
-}
-
-static void
-set_loc_shader_ptr(struct radv_shader_args *args, int idx, uint8_t *sgpr_idx)
-{
- bool use_32bit_pointers = idx != AC_UD_SCRATCH_RING_OFFSETS;
-
- set_loc_shader(args, idx, sgpr_idx, use_32bit_pointers ? 1 : 2);
-}
-
-static void
-set_loc_desc(struct radv_shader_args *args, int idx, uint8_t *sgpr_idx)
-{
- struct radv_userdata_locations *locs =
- &args->shader_info->user_sgprs_locs;
- struct radv_userdata_info *ud_info = &locs->descriptor_sets[idx];
- assert(ud_info);
-
- set_loc(ud_info, sgpr_idx, 1);
-
- locs->descriptor_sets_enabled |= 1 << idx;
-}
-
-struct user_sgpr_info {
- bool need_ring_offsets;
- bool indirect_all_descriptor_sets;
- uint8_t remaining_sgprs;
-};
-
-static bool needs_view_index_sgpr(struct radv_shader_args *args,
- gl_shader_stage stage)
-{
- switch (stage) {
- case MESA_SHADER_VERTEX:
- if (args->shader_info->needs_multiview_view_index ||
- (!args->options->key.vs_common_out.as_es && !args->options->key.vs_common_out.as_ls && args->options->key.has_multiview_view_index))
- return true;
- break;
- case MESA_SHADER_TESS_EVAL:
- if (args->shader_info->needs_multiview_view_index || (!args->options->key.vs_common_out.as_es && args->options->key.has_multiview_view_index))
- return true;
- break;
- case MESA_SHADER_GEOMETRY:
- case MESA_SHADER_TESS_CTRL:
- if (args->shader_info->needs_multiview_view_index)
- return true;
- break;
- default:
- break;
- }
- return false;
-}
-
-static uint8_t
-count_vs_user_sgprs(struct radv_shader_args *args)
-{
- uint8_t count = 0;
-
- if (args->shader_info->vs.has_vertex_buffers)
- count++;
- count += args->shader_info->vs.needs_draw_id ? 3 : 2;
-
- return count;
-}
-
-static void allocate_inline_push_consts(struct radv_shader_args *args,
- struct user_sgpr_info *user_sgpr_info)
-{
- uint8_t remaining_sgprs = user_sgpr_info->remaining_sgprs;
-
- /* Only supported if shaders use push constants. */
- if (args->shader_info->min_push_constant_used == UINT8_MAX)
- return;
-
- /* Only supported if shaders don't have indirect push constants. */
- if (args->shader_info->has_indirect_push_constants)
- return;
-
- /* Only supported for 32-bit push constants. */
- if (!args->shader_info->has_only_32bit_push_constants)
- return;
-
- uint8_t num_push_consts =
- (args->shader_info->max_push_constant_used -
- args->shader_info->min_push_constant_used) / 4;
-
- /* Check if the number of user SGPRs is large enough. */
- if (num_push_consts < remaining_sgprs) {
- args->shader_info->num_inline_push_consts = num_push_consts;
- } else {
- args->shader_info->num_inline_push_consts = remaining_sgprs;
- }
-
- /* Clamp to the maximum number of allowed inlined push constants. */
- if (args->shader_info->num_inline_push_consts > AC_MAX_INLINE_PUSH_CONSTS)
- args->shader_info->num_inline_push_consts = AC_MAX_INLINE_PUSH_CONSTS;
-
- if (args->shader_info->num_inline_push_consts == num_push_consts &&
- !args->shader_info->loads_dynamic_offsets) {
- /* Disable the default push constants path if all constants are
- * inlined and if shaders don't use dynamic descriptors.
- */
- args->shader_info->loads_push_constants = false;
- }
-
- args->shader_info->base_inline_push_consts =
- args->shader_info->min_push_constant_used / 4;
-}
-
-static void allocate_user_sgprs(struct radv_shader_args *args,
- gl_shader_stage stage,
- bool has_previous_stage,
- gl_shader_stage previous_stage,
- bool needs_view_index,
- struct user_sgpr_info *user_sgpr_info)
-{
- uint8_t user_sgpr_count = 0;
-
- memset(user_sgpr_info, 0, sizeof(struct user_sgpr_info));
-
- /* until we sort out scratch/global buffers always assign ring offsets for gs/vs/es */
- if (stage == MESA_SHADER_GEOMETRY ||
- stage == MESA_SHADER_VERTEX ||
- stage == MESA_SHADER_TESS_CTRL ||
- stage == MESA_SHADER_TESS_EVAL ||
- args->is_gs_copy_shader)
- user_sgpr_info->need_ring_offsets = true;
-
- if (stage == MESA_SHADER_FRAGMENT &&
- args->shader_info->ps.needs_sample_positions)
- user_sgpr_info->need_ring_offsets = true;
-
- /* 2 user sgprs will nearly always be allocated for scratch/rings */
- if (args->options->supports_spill || user_sgpr_info->need_ring_offsets) {
- user_sgpr_count += 2;
- }
-
- switch (stage) {
- case MESA_SHADER_COMPUTE:
- if (args->shader_info->cs.uses_grid_size)
- user_sgpr_count += 3;
- break;
- case MESA_SHADER_FRAGMENT:
- user_sgpr_count += args->shader_info->ps.needs_sample_positions;
- break;
- case MESA_SHADER_VERTEX:
- if (!args->is_gs_copy_shader)
- user_sgpr_count += count_vs_user_sgprs(args);
- break;
- case MESA_SHADER_TESS_CTRL:
- if (has_previous_stage) {
- if (previous_stage == MESA_SHADER_VERTEX)
- user_sgpr_count += count_vs_user_sgprs(args);
- }
- break;
- case MESA_SHADER_TESS_EVAL:
- break;
- case MESA_SHADER_GEOMETRY:
- if (has_previous_stage) {
- if (previous_stage == MESA_SHADER_VERTEX) {
- user_sgpr_count += count_vs_user_sgprs(args);
- }
- }
- break;
- default:
- break;
- }
-
- if (needs_view_index)
- user_sgpr_count++;
-
- if (args->shader_info->loads_push_constants)
- user_sgpr_count++;
-
- if (args->shader_info->so.num_outputs)
- user_sgpr_count++;
-
- uint32_t available_sgprs = args->options->chip_class >= GFX9 && stage != MESA_SHADER_COMPUTE ? 32 : 16;
- uint32_t remaining_sgprs = available_sgprs - user_sgpr_count;
- uint32_t num_desc_set =
- util_bitcount(args->shader_info->desc_set_used_mask);
-
- if (remaining_sgprs < num_desc_set) {
- user_sgpr_info->indirect_all_descriptor_sets = true;
- user_sgpr_info->remaining_sgprs = remaining_sgprs - 1;
- } else {
- user_sgpr_info->remaining_sgprs = remaining_sgprs - num_desc_set;
- }
-
- allocate_inline_push_consts(args, user_sgpr_info);
-}
-
-static void
-declare_global_input_sgprs(struct radv_shader_args *args,
- const struct user_sgpr_info *user_sgpr_info)
-{
- /* 1 for each descriptor set */
- if (!user_sgpr_info->indirect_all_descriptor_sets) {
- uint32_t mask = args->shader_info->desc_set_used_mask;
-
- while (mask) {
- int i = u_bit_scan(&mask);
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_CONST_PTR,
- &args->descriptor_sets[i]);
- }
- } else {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_CONST_PTR_PTR,
- &args->descriptor_sets[0]);
- }
-
- if (args->shader_info->loads_push_constants) {
- /* 1 for push constants and dynamic descriptors */
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_CONST_PTR,
- &args->ac.push_constants);
- }
-
- for (unsigned i = 0; i < args->shader_info->num_inline_push_consts; i++) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.inline_push_consts[i]);
- }
- args->ac.num_inline_push_consts = args->shader_info->num_inline_push_consts;
- args->ac.base_inline_push_consts = args->shader_info->base_inline_push_consts;
-
- if (args->shader_info->so.num_outputs) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR,
- &args->streamout_buffers);
- }
-}
-
-static void
-declare_vs_specific_input_sgprs(struct radv_shader_args *args,
- gl_shader_stage stage,
- bool has_previous_stage,
- gl_shader_stage previous_stage)
-{
- if (!args->is_gs_copy_shader &&
- (stage == MESA_SHADER_VERTEX ||
- (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
- if (args->shader_info->vs.has_vertex_buffers) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR,
- &args->vertex_buffers);
- }
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->ac.base_vertex);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->ac.start_instance);
- if (args->shader_info->vs.needs_draw_id) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->ac.draw_id);
- }
- }
-}
-
-static void
-declare_vs_input_vgprs(struct radv_shader_args *args)
-{
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.vertex_id);
- if (!args->is_gs_copy_shader) {
- if (args->options->key.vs_common_out.as_ls) {
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->rel_auto_id);
- if (args->options->chip_class >= GFX10) {
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* user vgpr */
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.instance_id);
- } else {
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.instance_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* unused */
- }
- } else {
- if (args->options->chip_class >= GFX10) {
- if (args->options->key.vs_common_out.as_ngg) {
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* user vgpr */
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* user vgpr */
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.instance_id);
- } else {
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* unused */
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->vs_prim_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.instance_id);
- }
- } else {
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.instance_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->vs_prim_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* unused */
- }
- }
- }
-}
-
-static void
-declare_streamout_sgprs(struct radv_shader_args *args, gl_shader_stage stage)
-{
- int i;
-
- if (args->options->use_ngg_streamout) {
- if (stage == MESA_SHADER_TESS_EVAL)
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL);
- return;
- }
-
- /* Streamout SGPRs. */
- if (args->shader_info->so.num_outputs) {
- assert(stage == MESA_SHADER_VERTEX ||
- stage == MESA_SHADER_TESS_EVAL);
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->streamout_config);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->streamout_write_idx);
- } else if (stage == MESA_SHADER_TESS_EVAL) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL);
- }
-
- /* A streamout buffer offset is loaded if the stride is non-zero. */
- for (i = 0; i < 4; i++) {
- if (!args->shader_info->so.strides[i])
- continue;
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->streamout_offset[i]);
- }
-}
-
-static void
-declare_tes_input_vgprs(struct radv_shader_args *args)
-{
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &args->tes_u);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &args->tes_v);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->tes_rel_patch_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.tes_patch_id);
-}
-
-static void
-set_global_input_locs(struct radv_shader_args *args,
- const struct user_sgpr_info *user_sgpr_info,
- uint8_t *user_sgpr_idx)
-{
- uint32_t mask = args->shader_info->desc_set_used_mask;
-
- if (!user_sgpr_info->indirect_all_descriptor_sets) {
- while (mask) {
- int i = u_bit_scan(&mask);
-
- set_loc_desc(args, i, user_sgpr_idx);
- }
- } else {
- set_loc_shader_ptr(args, AC_UD_INDIRECT_DESCRIPTOR_SETS,
- user_sgpr_idx);
-
- args->shader_info->need_indirect_descriptor_sets = true;
- }
-
- if (args->shader_info->loads_push_constants) {
- set_loc_shader_ptr(args, AC_UD_PUSH_CONSTANTS, user_sgpr_idx);
- }
-
- if (args->shader_info->num_inline_push_consts) {
- set_loc_shader(args, AC_UD_INLINE_PUSH_CONSTANTS, user_sgpr_idx,
- args->shader_info->num_inline_push_consts);
- }
-
- if (args->streamout_buffers.used) {
- set_loc_shader_ptr(args, AC_UD_STREAMOUT_BUFFERS,
- user_sgpr_idx);
- }
-}
-
static void
load_descriptor_sets(struct radv_shader_context *ctx)
{
}
}
-
-static void
-set_vs_specific_input_locs(struct radv_shader_args *args,
- gl_shader_stage stage, bool has_previous_stage,
- gl_shader_stage previous_stage,
- uint8_t *user_sgpr_idx)
-{
- if (!args->is_gs_copy_shader &&
- (stage == MESA_SHADER_VERTEX ||
- (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
- if (args->shader_info->vs.has_vertex_buffers) {
- set_loc_shader_ptr(args, AC_UD_VS_VERTEX_BUFFERS,
- user_sgpr_idx);
- }
-
- unsigned vs_num = 2;
- if (args->shader_info->vs.needs_draw_id)
- vs_num++;
-
- set_loc_shader(args, AC_UD_VS_BASE_VERTEX_START_INSTANCE,
- user_sgpr_idx, vs_num);
- }
-}
-
static enum ac_llvm_calling_convention
get_llvm_calling_convention(LLVMValueRef func, gl_shader_stage stage)
{
return stage == MESA_SHADER_VERTEX || stage == MESA_SHADER_TESS_EVAL;
}
-static void declare_inputs(struct radv_shader_args *args,
- gl_shader_stage stage,
- bool has_previous_stage,
- gl_shader_stage previous_stage)
-{
- struct user_sgpr_info user_sgpr_info;
- bool needs_view_index = needs_view_index_sgpr(args, stage);
-
- if (args->options->chip_class >= GFX10) {
- if (is_pre_gs_stage(stage) && args->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;
- }
- }
-
- for (int i = 0; i < MAX_SETS; i++)
- args->shader_info->user_sgprs_locs.descriptor_sets[i].sgpr_idx = -1;
- for (int i = 0; i < AC_UD_MAX_UD; i++)
- args->shader_info->user_sgprs_locs.shader_data[i].sgpr_idx = -1;
-
-
- allocate_user_sgprs(args, stage, has_previous_stage,
- previous_stage, needs_view_index, &user_sgpr_info);
-
- if (user_sgpr_info.need_ring_offsets && !args->options->supports_spill) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_DESC_PTR,
- &args->ring_offsets);
- }
-
- switch (stage) {
- case MESA_SHADER_COMPUTE:
- declare_global_input_sgprs(args, &user_sgpr_info);
-
- if (args->shader_info->cs.uses_grid_size) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 3, AC_ARG_INT,
- &args->ac.num_work_groups);
- }
-
- for (int i = 0; i < 3; i++) {
- if (args->shader_info->cs.uses_block_id[i]) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.workgroup_ids[i]);
- }
- }
-
- if (args->shader_info->cs.uses_local_invocation_idx) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.tg_size);
- }
-
- ac_add_arg(&args->ac, AC_ARG_VGPR, 3, AC_ARG_INT,
- &args->ac.local_invocation_ids);
- break;
- case MESA_SHADER_VERTEX:
- declare_global_input_sgprs(args, &user_sgpr_info);
-
- declare_vs_specific_input_sgprs(args, stage, has_previous_stage,
- previous_stage);
-
- if (needs_view_index) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.view_index);
- }
-
- if (args->options->key.vs_common_out.as_es) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->es2gs_offset);
- } else if (args->options->key.vs_common_out.as_ls) {
- /* no extra parameters */
- } else {
- declare_streamout_sgprs(args, stage);
- }
-
- declare_vs_input_vgprs(args);
- break;
- case MESA_SHADER_TESS_CTRL:
- if (has_previous_stage) {
- // First 6 system regs
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->oc_lds);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->merged_wave_info);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->tess_factor_offset);
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); // scratch offset
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); // unknown
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); // unknown
-
- declare_global_input_sgprs(args, &user_sgpr_info);
-
- declare_vs_specific_input_sgprs(args, stage,
- has_previous_stage,
- previous_stage);
-
- if (needs_view_index) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.view_index);
- }
-
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->ac.tcs_patch_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->ac.tcs_rel_ids);
-
- declare_vs_input_vgprs(args);
- } else {
- declare_global_input_sgprs(args, &user_sgpr_info);
-
- if (needs_view_index) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.view_index);
- }
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->oc_lds);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->tess_factor_offset);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->ac.tcs_patch_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->ac.tcs_rel_ids);
- }
- break;
- case MESA_SHADER_TESS_EVAL:
- declare_global_input_sgprs(args, &user_sgpr_info);
-
- if (needs_view_index)
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.view_index);
-
- if (args->options->key.vs_common_out.as_es) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->oc_lds);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->es2gs_offset);
- } else {
- declare_streamout_sgprs(args, stage);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->oc_lds);
- }
- declare_tes_input_vgprs(args);
- break;
- case MESA_SHADER_GEOMETRY:
- if (has_previous_stage) {
- // First 6 system regs
- if (args->options->key.vs_common_out.as_ngg) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->gs_tg_info);
- } else {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->gs2vs_offset);
- }
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->merged_wave_info);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->oc_lds);
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); // scratch offset
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); // unknown
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); // unknown
-
- declare_global_input_sgprs(args, &user_sgpr_info);
-
- if (previous_stage != MESA_SHADER_TESS_EVAL) {
- declare_vs_specific_input_sgprs(args, stage,
- has_previous_stage,
- previous_stage);
- }
-
- if (needs_view_index) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.view_index);
- }
-
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[0]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[2]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->ac.gs_prim_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->ac.gs_invocation_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[4]);
-
- if (previous_stage == MESA_SHADER_VERTEX) {
- declare_vs_input_vgprs(args);
- } else {
- declare_tes_input_vgprs(args);
- }
- } else {
- declare_global_input_sgprs(args, &user_sgpr_info);
-
- if (needs_view_index) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.view_index);
- }
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->gs2vs_offset);
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->gs_wave_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[0]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[1]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->ac.gs_prim_id);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[2]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[3]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[4]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->gs_vtx_offset[5]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT,
- &args->ac.gs_invocation_id);
- }
- break;
- case MESA_SHADER_FRAGMENT:
- declare_global_input_sgprs(args, &user_sgpr_info);
-
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->ac.prim_mask);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 2, AC_ARG_INT, &args->ac.persp_sample);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 2, AC_ARG_INT, &args->ac.persp_center);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 2, AC_ARG_INT, &args->ac.persp_centroid);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 3, AC_ARG_INT, NULL); /* persp pull model */
- ac_add_arg(&args->ac, AC_ARG_VGPR, 2, AC_ARG_INT, &args->ac.linear_sample);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 2, AC_ARG_INT, &args->ac.linear_center);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 2, AC_ARG_INT, &args->ac.linear_centroid);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_FLOAT, NULL); /* line stipple tex */
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &args->ac.frag_pos[0]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &args->ac.frag_pos[1]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &args->ac.frag_pos[2]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &args->ac.frag_pos[3]);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.front_face);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.ancillary);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.sample_coverage);
- ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* fixed pt */
- break;
- default:
- unreachable("Shader stage not implemented");
- }
-
- args->shader_info->num_input_vgprs = 0;
- args->shader_info->num_input_sgprs = args->options->supports_spill ? 2 : 0;
- args->shader_info->num_input_sgprs += args->ac.num_sgprs_used;
-
- if (stage != MESA_SHADER_FRAGMENT)
- args->shader_info->num_input_vgprs = args->ac.num_vgprs_used;
-
- uint8_t user_sgpr_idx = 0;
-
- if (args->options->supports_spill || user_sgpr_info.need_ring_offsets) {
- set_loc_shader_ptr(args, AC_UD_SCRATCH_RING_OFFSETS,
- &user_sgpr_idx);
- }
-
- /* 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(args, &user_sgpr_info, &user_sgpr_idx);
-
- switch (stage) {
- case MESA_SHADER_COMPUTE:
- if (args->shader_info->cs.uses_grid_size) {
- set_loc_shader(args, AC_UD_CS_GRID_SIZE,
- &user_sgpr_idx, 3);
- }
- break;
- case MESA_SHADER_VERTEX:
- set_vs_specific_input_locs(args, stage, has_previous_stage,
- previous_stage, &user_sgpr_idx);
- if (args->ac.view_index.used)
- set_loc_shader(args, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
- break;
- case MESA_SHADER_TESS_CTRL:
- set_vs_specific_input_locs(args, stage, has_previous_stage,
- previous_stage, &user_sgpr_idx);
- if (args->ac.view_index.used)
- set_loc_shader(args, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
- break;
- case MESA_SHADER_TESS_EVAL:
- if (args->ac.view_index.used)
- set_loc_shader(args, 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(args, stage,
- has_previous_stage,
- previous_stage,
- &user_sgpr_idx);
- }
- if (args->ac.view_index.used)
- set_loc_shader(args, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
- break;
- case MESA_SHADER_FRAGMENT:
- break;
- default:
- unreachable("Shader stage not implemented");
- }
-
- args->shader_info->num_user_sgprs = user_sgpr_idx;
-}
-
static void create_function(struct radv_shader_context *ctx,
gl_shader_stage stage,
bool has_previous_stage)
ctx->max_workgroup_size,
ctx->args->options);
- if (ctx->args->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), "");
- } else if (ctx->args->ring_offsets.used) {
- ctx->ring_offsets = ac_get_arg(&ctx->ac, ctx->args->ring_offsets);
- }
+ 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), "");
load_descriptor_sets(ctx);
if (ctx->ac.chip_class >= GFX10) {
desc_type |= S_008F0C_FORMAT(V_008F0C_IMG_FORMAT_32_FLOAT) |
- S_008F0C_OOB_SELECT(3) |
+ S_008F0C_OOB_SELECT(V_008F0C_OOB_SELECT_RAW) |
S_008F0C_RESOURCE_LEVEL(1);
} else {
desc_type |= S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) |
if (!is_tess_factor && writemask != 0xF)
ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, value, 1,
buf_addr, oc_lds,
- 4 * (base + chan), ac_glc, false);
+ 4 * (base + chan), ac_glc);
}
if (writemask == 0xF) {
ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, src, 4,
buf_addr, oc_lds,
- (base * 4), ac_glc, false);
+ (base * 4), ac_glc);
}
}
return result;
}
-
-static void radv_emit_kill(struct ac_shader_abi *abi, LLVMValueRef visible)
-{
- struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
- ac_build_kill_if_false(&ctx->ac, visible);
-}
-
static uint32_t
radv_get_sample_pos_offset(uint32_t num_samples)
{
static void gfx10_ngg_gs_emit_vertex(struct radv_shader_context *ctx,
unsigned stream,
+ LLVMValueRef vertexidx,
LLVMValueRef *addrs);
static void
-visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addrs)
+visit_emit_vertex_with_counter(struct ac_shader_abi *abi, unsigned stream,
+ LLVMValueRef vertexidx, LLVMValueRef *addrs)
{
- LLVMValueRef gs_next_vertex;
- LLVMValueRef can_emit;
unsigned offset = 0;
struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
if (ctx->args->options->key.vs_common_out.as_ngg) {
- gfx10_ngg_gs_emit_vertex(ctx, stream, addrs);
+ gfx10_ngg_gs_emit_vertex(ctx, stream, vertexidx, addrs);
return;
}
- /* Write vertex attribute values to GSVS ring */
- gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
- ctx->gs_next_vertex[stream],
- "");
-
- /* If this thread has already emitted the declared maximum number of
- * vertices, don't emit any more: excessive vertex emissions are not
- * supposed to have any effect.
- */
- can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex,
- LLVMConstInt(ctx->ac.i32, ctx->shader->info.gs.vertices_out, false), "");
-
- bool use_kill = !ctx->args->shader_info->gs.writes_memory;
- if (use_kill)
- ac_build_kill_if_false(&ctx->ac, can_emit);
- else
- ac_build_ifcc(&ctx->ac, can_emit, 6505);
-
for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
unsigned output_usage_mask =
ctx->args->shader_info->gs.output_usage_mask[i];
offset++;
- voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
+ voffset = LLVMBuildAdd(ctx->ac.builder, voffset, vertexidx, "");
voffset = LLVMBuildMul(ctx->ac.builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
out_val = ac_to_integer(&ctx->ac, out_val);
voffset,
ac_get_arg(&ctx->ac,
ctx->args->gs2vs_offset),
- 0, ac_glc | ac_slc, true);
+ 0, ac_glc | ac_slc | ac_swizzled);
}
}
- gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex,
- ctx->ac.i32_1, "");
- LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
-
ac_build_sendmsg(&ctx->ac,
AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
ctx->gs_wave_id);
-
- if (!use_kill)
- ac_build_endif(&ctx->ac, 6505);
}
static void
return LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.i32, "");
}
-static unsigned
-get_num_channels_from_data_format(unsigned data_format)
-{
- switch (data_format) {
- case V_008F0C_BUF_DATA_FORMAT_8:
- case V_008F0C_BUF_DATA_FORMAT_16:
- case V_008F0C_BUF_DATA_FORMAT_32:
- return 1;
- case V_008F0C_BUF_DATA_FORMAT_8_8:
- case V_008F0C_BUF_DATA_FORMAT_16_16:
- case V_008F0C_BUF_DATA_FORMAT_32_32:
- return 2;
- case V_008F0C_BUF_DATA_FORMAT_10_11_11:
- case V_008F0C_BUF_DATA_FORMAT_11_11_10:
- case V_008F0C_BUF_DATA_FORMAT_32_32_32:
- return 3;
- case V_008F0C_BUF_DATA_FORMAT_8_8_8_8:
- case V_008F0C_BUF_DATA_FORMAT_10_10_10_2:
- case V_008F0C_BUF_DATA_FORMAT_2_10_10_10:
- case V_008F0C_BUF_DATA_FORMAT_16_16_16_16:
- case V_008F0C_BUF_DATA_FORMAT_32_32_32_32:
- return 4;
- default:
- break;
- }
-
- return 4;
-}
-
static LLVMValueRef
radv_fixup_vertex_input_fetches(struct radv_shader_context *ctx,
LLVMValueRef value,
for (unsigned i = 0; i < num_channels; i++)
chan[i] = ac_llvm_extract_elem(&ctx->ac, value, i);
} else {
- if (num_channels) {
- assert(num_channels == 1);
- chan[0] = value;
- }
+ assert(num_channels == 1);
+ chan[0] = value;
}
for (unsigned i = num_channels; i < 4; i++) {
ctx->args->ac.base_vertex), "");
}
+ const struct ac_data_format_info *vtx_info = ac_get_data_format_info(data_format);
+
/* Adjust the number of channels to load based on the vertex
* attribute format.
*/
- unsigned num_format_channels = get_num_channels_from_data_format(data_format);
- unsigned num_channels = MIN2(num_input_channels, num_format_channels);
+ unsigned num_channels = MIN2(num_input_channels, vtx_info->num_channels);
unsigned attrib_binding = ctx->args->options->key.vs.vertex_attribute_bindings[attrib_index];
unsigned attrib_offset = ctx->args->options->key.vs.vertex_attribute_offsets[attrib_index];
unsigned attrib_stride = ctx->args->options->key.vs.vertex_attribute_strides[attrib_index];
num_channels = MAX2(num_channels, 3);
}
- if (attrib_stride != 0 && attrib_offset > attrib_stride) {
- LLVMValueRef buffer_offset =
- LLVMConstInt(ctx->ac.i32,
- attrib_offset / attrib_stride, false);
+ t_offset = LLVMConstInt(ctx->ac.i32, attrib_binding, false);
+ t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
- buffer_index = LLVMBuildAdd(ctx->ac.builder,
- buffer_index,
- buffer_offset, "");
+ /* Perform per-channel vertex fetch operations if unaligned
+ * access are detected. Only GFX6 and GFX10 are affected.
+ */
+ bool unaligned_vertex_fetches = false;
+ if ((ctx->ac.chip_class == GFX6 || ctx->ac.chip_class >= GFX10) &&
+ vtx_info->chan_format != data_format &&
+ ((attrib_offset % vtx_info->element_size) ||
+ (attrib_stride % vtx_info->element_size)))
+ unaligned_vertex_fetches = true;
+
+ if (unaligned_vertex_fetches) {
+ unsigned chan_format = vtx_info->chan_format;
+ LLVMValueRef values[4];
- attrib_offset = attrib_offset % attrib_stride;
- }
+ assert(ctx->ac.chip_class == GFX6 ||
+ ctx->ac.chip_class >= GFX10);
- t_offset = LLVMConstInt(ctx->ac.i32, attrib_binding, false);
- t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
+ for (unsigned chan = 0; chan < num_channels; chan++) {
+ unsigned chan_offset = attrib_offset + chan * vtx_info->chan_byte_size;
+ LLVMValueRef chan_index = buffer_index;
+
+ if (attrib_stride != 0 && chan_offset > attrib_stride) {
+ LLVMValueRef buffer_offset =
+ LLVMConstInt(ctx->ac.i32,
+ chan_offset / attrib_stride, false);
- input = ac_build_struct_tbuffer_load(&ctx->ac, t_list,
- buffer_index,
- LLVMConstInt(ctx->ac.i32, attrib_offset, false),
- ctx->ac.i32_0, ctx->ac.i32_0,
- num_channels,
- data_format, num_format, 0, true);
+ chan_index = LLVMBuildAdd(ctx->ac.builder,
+ buffer_index,
+ buffer_offset, "");
+
+ chan_offset = chan_offset % attrib_stride;
+ }
+
+ values[chan] = ac_build_struct_tbuffer_load(&ctx->ac, t_list,
+ chan_index,
+ LLVMConstInt(ctx->ac.i32, chan_offset, false),
+ ctx->ac.i32_0, ctx->ac.i32_0, 1,
+ chan_format, num_format, 0, true);
+ }
+
+ input = ac_build_gather_values(&ctx->ac, values, num_channels);
+ } else {
+ if (attrib_stride != 0 && attrib_offset > attrib_stride) {
+ LLVMValueRef buffer_offset =
+ LLVMConstInt(ctx->ac.i32,
+ attrib_offset / attrib_stride, false);
+
+ buffer_index = LLVMBuildAdd(ctx->ac.builder,
+ buffer_index,
+ buffer_offset, "");
+
+ attrib_offset = attrib_offset % attrib_stride;
+ }
+
+ input = ac_build_struct_tbuffer_load(&ctx->ac, t_list,
+ buffer_index,
+ LLVMConstInt(ctx->ac.i32, attrib_offset, false),
+ ctx->ac.i32_0, ctx->ac.i32_0,
+ num_channels,
+ data_format, num_format, 0, true);
+ }
if (ctx->args->options->key.vs.post_shuffle & (1 << attrib_index)) {
LLVMValueRef c[4];
break;
}
+ /* Replace NaN by zero (only 32-bit) to fix game bugs if
+ * requested.
+ */
+ if (ctx->args->options->enable_mrt_output_nan_fixup &&
+ !is_16bit &&
+ (col_format == V_028714_SPI_SHADER_32_R ||
+ col_format == V_028714_SPI_SHADER_32_GR ||
+ col_format == V_028714_SPI_SHADER_32_AR ||
+ col_format == V_028714_SPI_SHADER_32_ABGR ||
+ col_format == V_028714_SPI_SHADER_FP16_ABGR)) {
+ for (unsigned i = 0; i < 4; i++) {
+ LLVMValueRef args[2] = {
+ values[i],
+ LLVMConstInt(ctx->ac.i32, S_NAN | Q_NAN, false)
+ };
+ LLVMValueRef isnan =
+ ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.class.f32", ctx->ac.i1,
+ args, 2, AC_FUNC_ATTR_READNONE);
+ values[i] = LLVMBuildSelect(ctx->ac.builder, isnan,
+ ctx->ac.f32_0,
+ values[i], "");
+ }
+ }
+
/* Pack f16 or norm_i16/u16. */
if (packf) {
for (chan = 0; chan < 2; chan++) {
ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf],
vdata, num_comps, so_write_offsets[buf],
ctx->ac.i32_0, offset,
- ac_glc | ac_slc, false);
+ ac_glc | ac_slc);
}
static void
if (slot_name != VARYING_SLOT_LAYER &&
slot_name != VARYING_SLOT_PRIMITIVE_ID &&
+ slot_name != VARYING_SLOT_VIEWPORT &&
slot_name != VARYING_SLOT_CLIP_DIST0 &&
slot_name != VARYING_SLOT_CLIP_DIST1 &&
slot_name < VARYING_SLOT_VAR0)
outinfo->pos_exports++;
}
- /* Navi10-14 skip POS0 exports if EXEC=0 and DONE=0, causing a hang.
+ /* GFX10 skip POS0 exports if EXEC=0 and DONE=0, causing a hang.
* Setting valid_mask=1 prevents it and has no other effect.
*/
- if (ctx->ac.family == CHIP_NAVI10 ||
- ctx->ac.family == CHIP_NAVI12 ||
- ctx->ac.family == CHIP_NAVI14)
+ if (ctx->ac.chip_class == GFX10)
pos_args[0].valid_mask = 1;
pos_idx = 0;
NULL,
ac_get_arg(&ctx->ac, ctx->args->es2gs_offset),
(4 * param_index + j) * 4,
- ac_glc | ac_slc, true);
+ ac_glc | ac_slc | ac_swizzled);
}
}
}
{
return ac_build_bfe(&ctx->ac, ac_get_arg(&ctx->ac, ctx->args->gs_tg_info),
ctx->ac.i32_0,
- LLVMConstInt(ctx->ac.i32, 11, false),
+ LLVMConstInt(ctx->ac.i32, 12, false),
false);
}
return ngg_gs_vertex_ptr(ctx, vertexidx);
}
-/* Send GS Alloc Req message from the first wave of the group to SPI.
- * Message payload is:
- * - bits 0..10: vertices in group
- * - bits 12..22: primitives in group
- */
-static void build_sendmsg_gs_alloc_req(struct radv_shader_context *ctx,
- LLVMValueRef vtx_cnt,
- LLVMValueRef prim_cnt)
+static LLVMValueRef
+ngg_gs_get_emit_output_ptr(struct radv_shader_context *ctx, LLVMValueRef vertexptr,
+ unsigned out_idx)
{
- LLVMBuilderRef builder = ctx->ac.builder;
- LLVMValueRef tmp;
-
- tmp = LLVMBuildICmp(builder, LLVMIntEQ, get_wave_id_in_tg(ctx), ctx->ac.i32_0, "");
- ac_build_ifcc(&ctx->ac, tmp, 5020);
-
- tmp = LLVMBuildShl(builder, prim_cnt, LLVMConstInt(ctx->ac.i32, 12, false),"");
- tmp = LLVMBuildOr(builder, tmp, vtx_cnt, "");
- ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_ALLOC_REQ, tmp);
-
- ac_build_endif(&ctx->ac, 5020);
+ LLVMValueRef gep_idx[3] = {
+ ctx->ac.i32_0, /* implied C-style array */
+ ctx->ac.i32_0, /* first struct entry */
+ LLVMConstInt(ctx->ac.i32, out_idx, false),
+ };
+ return LLVMBuildGEP(ctx->ac.builder, vertexptr, gep_idx, 3, "");
}
-struct ngg_prim {
- unsigned num_vertices;
- LLVMValueRef isnull;
- LLVMValueRef swap;
- LLVMValueRef index[3];
- LLVMValueRef edgeflag[3];
-};
-
-static void build_export_prim(struct radv_shader_context *ctx,
- const struct ngg_prim *prim)
+static LLVMValueRef
+ngg_gs_get_emit_primflag_ptr(struct radv_shader_context *ctx, LLVMValueRef vertexptr,
+ unsigned stream)
{
- LLVMBuilderRef builder = ctx->ac.builder;
- struct ac_export_args args;
- LLVMValueRef vertices[3];
- LLVMValueRef odd, even;
- LLVMValueRef tmp;
-
- tmp = LLVMBuildZExt(builder, prim->isnull, ctx->ac.i32, "");
- args.out[0] = LLVMBuildShl(builder, tmp, LLVMConstInt(ctx->ac.i32, 31, false), "");
-
- for (unsigned i = 0; i < prim->num_vertices; ++i) {
- tmp = LLVMBuildZExt(builder, prim->edgeflag[i], ctx->ac.i32, "");
- tmp = LLVMBuildShl(builder, tmp,
- LLVMConstInt(ctx->ac.i32, 9, false), "");
- vertices[i] = LLVMBuildOr(builder, prim->index[i], tmp, "");
- }
-
- switch (prim->num_vertices) {
- case 1:
- args.out[0] = LLVMBuildOr(builder, args.out[0], vertices[0], "");
- break;
- case 2:
- tmp = LLVMBuildShl(builder, vertices[1],
- LLVMConstInt(ctx->ac.i32, 10, false), "");
- tmp = LLVMBuildOr(builder, args.out[0], tmp, "");
- args.out[0] = LLVMBuildOr(builder, tmp, vertices[0], "");
- break;
- case 3:
- /* Swap vertices if needed to follow drawing order. */
- tmp = LLVMBuildShl(builder, vertices[2],
- LLVMConstInt(ctx->ac.i32, 20, false), "");
- even = LLVMBuildOr(builder, args.out[0], tmp, "");
- tmp = LLVMBuildShl(builder, vertices[1],
- LLVMConstInt(ctx->ac.i32, 10, false), "");
- even = LLVMBuildOr(builder, even, tmp, "");
- even = LLVMBuildOr(builder, even, vertices[0], "");
-
- tmp = LLVMBuildShl(builder, vertices[1],
- LLVMConstInt(ctx->ac.i32, 20, false), "");
- odd = LLVMBuildOr(builder, args.out[0], tmp, "");
- tmp = LLVMBuildShl(builder, vertices[2],
- LLVMConstInt(ctx->ac.i32, 10, false), "");
- odd = LLVMBuildOr(builder, odd, tmp, "");
- odd = LLVMBuildOr(builder, odd, vertices[0], "");
-
- args.out[0] = LLVMBuildSelect(builder, prim->swap, odd, even, "");
- break;
- default:
- unreachable("invalid number of vertices");
- }
-
- args.out[0] = LLVMBuildBitCast(builder, args.out[0], ctx->ac.f32, "");
- args.out[1] = LLVMGetUndef(ctx->ac.f32);
- args.out[2] = LLVMGetUndef(ctx->ac.f32);
- args.out[3] = LLVMGetUndef(ctx->ac.f32);
-
- args.target = V_008DFC_SQ_EXP_PRIM;
- args.enabled_channels = 1;
- args.done = true;
- args.valid_mask = false;
- args.compr = false;
-
- ac_build_export(&ctx->ac, &args);
+ LLVMValueRef gep_idx[3] = {
+ ctx->ac.i32_0, /* implied C-style array */
+ ctx->ac.i32_1, /* second struct entry */
+ LLVMConstInt(ctx->ac.i32, stream, false),
+ };
+ return LLVMBuildGEP(ctx->ac.builder, vertexptr, gep_idx, 3, "");
}
static struct radv_stream_output *
for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
unsigned output_usage_mask =
ctx->args->shader_info->gs.output_usage_mask[i];
- uint8_t output_stream =
- output_stream = ctx->args->shader_info->gs.output_streams[i];
+ uint8_t output_stream = ctx->args->shader_info->gs.output_streams[i];
if (!(ctx->output_mask & (1ull << i)) ||
output_stream != stream)
/* TODO: primitive culling */
- build_sendmsg_gs_alloc_req(ctx, ngg_get_vtx_cnt(ctx), ngg_get_prim_cnt(ctx));
+ ac_build_sendmsg_gs_alloc_req(&ctx->ac, get_wave_id_in_tg(ctx),
+ ngg_get_vtx_cnt(ctx), ngg_get_prim_cnt(ctx));
/* TODO: streamout queries */
- /* Export primitive data to the index buffer. Format is:
- * - bits 0..8: index 0
- * - bit 9: edge flag 0
- * - bits 10..18: index 1
- * - bit 19: edge flag 1
- * - bits 20..28: index 2
- * - bit 29: edge flag 2
- * - bit 31: null primitive (skip)
+ /* Export primitive data to the index buffer.
*
* For the first version, we will always build up all three indices
* independent of the primitive type. The additional garbage data
*/
ac_build_ifcc(&ctx->ac, is_gs_thread, 6001);
{
- struct ngg_prim prim = {};
-
- prim.num_vertices = num_vertices;
- prim.isnull = ctx->ac.i1false;
- prim.swap = ctx->ac.i1false;
- memcpy(prim.index, vtxindex, sizeof(vtxindex[0]) * 3);
+ struct ac_ngg_prim prim = {};
- for (unsigned i = 0; i < num_vertices; ++i) {
- tmp = LLVMBuildLShr(builder,
- ac_get_arg(&ctx->ac, ctx->args->ac.gs_invocation_id),
- LLVMConstInt(ctx->ac.i32, 8 + i, false), "");
- prim.edgeflag[i] = LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
+ if (ctx->args->options->key.vs_common_out.as_ngg_passthrough) {
+ prim.passthrough = ac_get_arg(&ctx->ac, ctx->args->gs_vtx_offset[0]);
+ } else {
+ prim.num_vertices = num_vertices;
+ prim.isnull = ctx->ac.i1false;
+ memcpy(prim.index, vtxindex, sizeof(vtxindex[0]) * 3);
+
+ for (unsigned i = 0; i < num_vertices; ++i) {
+ tmp = LLVMBuildLShr(builder,
+ ac_get_arg(&ctx->ac, ctx->args->ac.gs_invocation_id),
+ LLVMConstInt(ctx->ac.i32, 8 + i, false), "");
+ prim.edgeflag[i] = LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
+ }
}
- build_export_prim(ctx, &prim);
+ ac_build_export_prim(&ctx->ac, &prim);
}
ac_build_endif(&ctx->ac, 6001);
LLVMBuildStore(builder, tmp, ctx->gs_next_vertex[stream]);
tmp = ngg_gs_emit_vertex_ptr(ctx, gsthread, vertexidx);
- LLVMValueRef gep_idx[3] = {
- ctx->ac.i32_0, /* implied C-style array */
- ctx->ac.i32_1, /* second entry of struct */
- LLVMConstInt(ctx->ac.i32, stream, false),
- };
- tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
- LLVMBuildStore(builder, i8_0, tmp);
+ LLVMBuildStore(builder, i8_0,
+ ngg_gs_get_emit_primflag_ptr(ctx, tmp, stream));
ac_build_endloop(&ctx->ac, 5100);
}
if (!ctx->args->shader_info->gs.num_stream_output_components[stream])
continue;
- LLVMValueRef gep_idx[3] = {
- ctx->ac.i32_0, /* implicit C-style array */
- ctx->ac.i32_1, /* second value of struct */
- LLVMConstInt(ctx->ac.i32, stream, false),
- };
- tmp = LLVMBuildGEP(builder, vertexptr, gep_idx, 3, "");
- tmp = LLVMBuildLoad(builder, tmp, "");
+ tmp = LLVMBuildLoad(builder,
+ ngg_gs_get_emit_primflag_ptr(ctx, vertexptr, stream), "");
tmp = LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
tmp2 = LLVMBuildICmp(builder, LLVMIntULT, tid, num_emit_threads, "");
nggso.prim_enable[stream] = LLVMBuildAnd(builder, tmp, tmp2, "");
build_streamout(ctx, &nggso);
}
+ /* Write shader query data. */
+ tmp = ac_get_arg(&ctx->ac, ctx->args->ngg_gs_state);
+ tmp = LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
+ ac_build_ifcc(&ctx->ac, tmp, 5109);
+ tmp = LLVMBuildICmp(builder, LLVMIntULT, tid,
+ LLVMConstInt(ctx->ac.i32, 4, false), "");
+ ac_build_ifcc(&ctx->ac, tmp, 5110);
+ {
+ tmp = LLVMBuildLoad(builder, ac_build_gep0(&ctx->ac, ctx->gs_ngg_scratch, tid), "");
+
+ ac_llvm_add_target_dep_function_attr(ctx->main_function,
+ "amdgpu-gds-size", 256);
+
+ LLVMTypeRef gdsptr = LLVMPointerType(ctx->ac.i32, AC_ADDR_SPACE_GDS);
+ LLVMValueRef gdsbase = LLVMBuildIntToPtr(builder, ctx->ac.i32_0, gdsptr, "");
+
+ const char *sync_scope = LLVM_VERSION_MAJOR >= 9 ? "workgroup-one-as" : "workgroup";
+
+ /* Use a plain GDS atomic to accumulate the number of generated
+ * primitives.
+ */
+ ac_build_atomic_rmw(&ctx->ac, LLVMAtomicRMWBinOpAdd, gdsbase,
+ tmp, sync_scope);
+ }
+ ac_build_endif(&ctx->ac, 5110);
+ ac_build_endif(&ctx->ac, 5109);
+
/* TODO: culling */
/* Determine vertex liveness. */
/* Load primitive liveness */
tmp = ngg_gs_vertex_ptr(ctx, primidx);
- LLVMValueRef gep_idx[3] = {
- ctx->ac.i32_0, /* implicit C-style array */
- ctx->ac.i32_1, /* second value of struct */
- ctx->ac.i32_0, /* stream 0 */
- };
- tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
- tmp = LLVMBuildLoad(builder, tmp, "");
+ tmp = LLVMBuildLoad(builder,
+ ngg_gs_get_emit_primflag_ptr(ctx, tmp, 0), "");
const LLVMValueRef primlive =
LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
* there are 4 or more contiguous null primitives in the export
* (in the common case of single-dword prim exports).
*/
- build_sendmsg_gs_alloc_req(ctx, vertlive_scan.result_reduce, num_emit_threads);
+ ac_build_sendmsg_gs_alloc_req(&ctx->ac, get_wave_id_in_tg(ctx),
+ vertlive_scan.result_reduce, num_emit_threads);
/* Setup the reverse vertex compaction permutation. We re-use stream 1
* of the primitive liveness flags, relying on the fact that each
ac_build_ifcc(&ctx->ac, vertlive, 5130);
{
tmp = ngg_gs_vertex_ptr(ctx, vertlive_scan.result_exclusive);
- LLVMValueRef gep_idx[3] = {
- ctx->ac.i32_0, /* implicit C-style array */
- ctx->ac.i32_1, /* second value of struct */
- ctx->ac.i32_1, /* stream 1 */
- };
- tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
tmp2 = LLVMBuildTrunc(builder, tid, ctx->ac.i8, "");
- LLVMBuildStore(builder, tmp2, tmp);
+ LLVMBuildStore(builder, tmp2,
+ ngg_gs_get_emit_primflag_ptr(ctx, tmp, 1));
}
ac_build_endif(&ctx->ac, 5130);
tmp = LLVMBuildICmp(builder, LLVMIntULT, tid, num_emit_threads, "");
ac_build_ifcc(&ctx->ac, tmp, 5140);
{
- struct ngg_prim prim = {};
+ LLVMValueRef flags;
+ struct ac_ngg_prim prim = {};
prim.num_vertices = verts_per_prim;
tmp = ngg_gs_vertex_ptr(ctx, tid);
- LLVMValueRef gep_idx[3] = {
- ctx->ac.i32_0, /* implicit C-style array */
- ctx->ac.i32_1, /* second value of struct */
- ctx->ac.i32_0, /* primflag */
- };
- tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
- tmp = LLVMBuildLoad(builder, tmp, "");
- prim.isnull = LLVMBuildICmp(builder, LLVMIntEQ, tmp,
- LLVMConstInt(ctx->ac.i8, 0, false), "");
- prim.swap = LLVMBuildICmp(builder, LLVMIntEQ,
- LLVMBuildAnd(builder, tid, LLVMConstInt(ctx->ac.i32, 1, false), ""),
- LLVMConstInt(ctx->ac.i32, 1, false), "");
+ flags = LLVMBuildLoad(builder,
+ ngg_gs_get_emit_primflag_ptr(ctx, tmp, 0), "");
+ prim.isnull = LLVMBuildNot(builder, LLVMBuildTrunc(builder, flags, ctx->ac.i1, ""), "");
for (unsigned i = 0; i < verts_per_prim; ++i) {
prim.index[i] = LLVMBuildSub(builder, vertlive_scan.result_exclusive,
prim.edgeflag[i] = ctx->ac.i1false;
}
- build_export_prim(ctx, &prim);
+ /* Geometry shaders output triangle strips, but NGG expects
+ * triangles. We need to change the vertex order for odd
+ * triangles to get correct front/back facing by swapping 2
+ * vertex indices, but we also have to keep the provoking
+ * vertex in the same place.
+ */
+ if (verts_per_prim == 3) {
+ LLVMValueRef is_odd = LLVMBuildLShr(builder, flags, ctx->ac.i8_1, "");
+ is_odd = LLVMBuildTrunc(builder, is_odd, ctx->ac.i1, "");
+
+ struct ac_ngg_prim in = prim;
+ prim.index[0] = in.index[0];
+ prim.index[1] = LLVMBuildSelect(builder, is_odd,
+ in.index[2], in.index[1], "");
+ prim.index[2] = LLVMBuildSelect(builder, is_odd,
+ in.index[1], in.index[2], "");
+ }
+
+ ac_build_export_prim(&ctx->ac, &prim);
}
ac_build_endif(&ctx->ac, 5140);
outinfo->pos_exports = 0;
tmp = ngg_gs_vertex_ptr(ctx, tid);
- LLVMValueRef gep_idx[3] = {
- ctx->ac.i32_0, /* implicit C-style array */
- ctx->ac.i32_1, /* second value of struct */
- ctx->ac.i32_1, /* stream 1: source data index */
- };
- tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
- tmp = LLVMBuildLoad(builder, tmp, "");
+ tmp = LLVMBuildLoad(builder,
+ ngg_gs_get_emit_primflag_ptr(ctx, tmp, 1), "");
tmp = LLVMBuildZExt(builder, tmp, ctx->ac.i32, "");
const LLVMValueRef vertexptr = ngg_gs_vertex_ptr(ctx, tmp);
unsigned out_idx = 0;
- gep_idx[1] = ctx->ac.i32_0;
for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
unsigned output_usage_mask =
ctx->args->shader_info->gs.output_usage_mask[i];
if (!(output_usage_mask & (1 << j)))
continue;
- gep_idx[2] = LLVMConstInt(ctx->ac.i32, out_idx, false);
- tmp = LLVMBuildGEP(builder, vertexptr, gep_idx, 3, "");
+ tmp = ngg_gs_get_emit_output_ptr(ctx, vertexptr, out_idx);
tmp = LLVMBuildLoad(builder, tmp, "");
LLVMTypeRef type = LLVMGetAllocatedType(ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
static void gfx10_ngg_gs_emit_vertex(struct radv_shader_context *ctx,
unsigned stream,
+ LLVMValueRef vertexidx,
LLVMValueRef *addrs)
{
LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef tmp;
- const LLVMValueRef vertexidx =
- LLVMBuildLoad(builder, ctx->gs_next_vertex[stream], "");
-
- /* If this thread has already emitted the declared maximum number of
- * vertices, skip the write: excessive vertex emissions are not
- * supposed to have any effect.
- */
- const LLVMValueRef can_emit =
- LLVMBuildICmp(builder, LLVMIntULT, vertexidx,
- LLVMConstInt(ctx->ac.i32, ctx->shader->info.gs.vertices_out, false), "");
- ac_build_ifcc(&ctx->ac, can_emit, 9001);
-
- tmp = LLVMBuildAdd(builder, vertexidx, ctx->ac.i32_1, "");
- tmp = LLVMBuildSelect(builder, can_emit, tmp, vertexidx, "");
- LLVMBuildStore(builder, tmp, ctx->gs_next_vertex[stream]);
const LLVMValueRef vertexptr =
ngg_gs_emit_vertex_ptr(ctx, get_thread_id_in_tg(ctx), vertexidx);
LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder,
out_ptr[j], "");
- LLVMValueRef gep_idx[3] = {
- ctx->ac.i32_0, /* implied C-style array */
- ctx->ac.i32_0, /* first entry of struct */
- LLVMConstInt(ctx->ac.i32, out_idx, false),
- };
- LLVMValueRef ptr = LLVMBuildGEP(builder, vertexptr, gep_idx, 3, "");
-
out_val = ac_to_integer(&ctx->ac, out_val);
out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
- LLVMBuildStore(builder, out_val, ptr);
+ LLVMBuildStore(builder, out_val,
+ ngg_gs_get_emit_output_ptr(ctx, vertexptr, out_idx));
}
}
assert(out_idx * 4 <= ctx->args->shader_info->gs.gsvs_vertex_size);
+ /* Store the current number of emitted vertices to zero out remaining
+ * primitive flags in case the geometry shader doesn't emit the maximum
+ * number of vertices.
+ */
+ tmp = LLVMBuildAdd(builder, vertexidx, ctx->ac.i32_1, "");
+ LLVMBuildStore(builder, tmp, ctx->gs_next_vertex[stream]);
+
/* Determine and store whether this vertex completed a primitive. */
const LLVMValueRef curverts = LLVMBuildLoad(builder, ctx->gs_curprim_verts[stream], "");
const LLVMValueRef iscompleteprim =
LLVMBuildICmp(builder, LLVMIntUGE, curverts, tmp, "");
+ /* Since the geometry shader emits triangle strips, we need to
+ * track which primitive is odd and swap vertex indices to get
+ * the correct vertex order.
+ */
+ LLVMValueRef is_odd = ctx->ac.i1false;
+ if (stream == 0 &&
+ si_conv_gl_prim_to_vertices(ctx->shader->info.gs.output_primitive) == 3) {
+ tmp = LLVMBuildAnd(builder, curverts, ctx->ac.i32_1, "");
+ is_odd = LLVMBuildICmp(builder, LLVMIntEQ, tmp, ctx->ac.i32_1, "");
+ }
+
tmp = LLVMBuildAdd(builder, curverts, ctx->ac.i32_1, "");
LLVMBuildStore(builder, tmp, ctx->gs_curprim_verts[stream]);
- LLVMValueRef gep_idx[3] = {
- ctx->ac.i32_0, /* implied C-style array */
- ctx->ac.i32_1, /* second struct entry */
- LLVMConstInt(ctx->ac.i32, stream, false),
- };
- const LLVMValueRef primflagptr =
- LLVMBuildGEP(builder, vertexptr, gep_idx, 3, "");
-
+ /* The per-vertex primitive flag encoding:
+ * bit 0: whether this vertex finishes a primitive
+ * bit 1: whether the primitive is odd (if we are emitting triangle strips)
+ */
tmp = LLVMBuildZExt(builder, iscompleteprim, ctx->ac.i8, "");
- LLVMBuildStore(builder, tmp, primflagptr);
+ tmp = LLVMBuildOr(builder, tmp,
+ LLVMBuildShl(builder,
+ LLVMBuildZExt(builder, is_odd, ctx->ac.i8, ""),
+ ctx->ac.i8_1, ""), "");
+ LLVMBuildStore(builder, tmp,
+ ngg_gs_get_emit_primflag_ptr(ctx, vertexptr, stream));
tmp = LLVMBuildLoad(builder, ctx->gs_generated_prims[stream], "");
tmp = LLVMBuildAdd(builder, tmp, LLVMBuildZExt(builder, iscompleteprim, ctx->ac.i32, ""), "");
LLVMBuildStore(builder, tmp, ctx->gs_generated_prims[stream]);
-
- ac_build_endif(&ctx->ac, 9001);
}
static void
ac_build_buffer_store_dword(&ctx->ac, buffer,
LLVMConstInt(ctx->ac.i32, 0x80000000, false),
1, ctx->ac.i32_0, tf_base,
- 0, ac_glc, false);
+ 0, ac_glc);
tf_offset += 4;
ac_build_endif(&ctx->ac, 6504);
/* Store the tessellation factors. */
ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
MIN2(stride, 4), byteoffset, tf_base,
- tf_offset, ac_glc, false);
+ tf_offset, ac_glc);
if (vec1)
ac_build_buffer_store_dword(&ctx->ac, buffer, vec1,
stride - 4, byteoffset, tf_base,
- 16 + tf_offset, ac_glc, false);
+ 16 + tf_offset, ac_glc);
//store to offchip for TES to read - only if TES reads them
if (ctx->args->options->key.tcs.tes_reads_tess_factors) {
ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, outer_vec,
outer_comps, tf_outer_offset,
ac_get_arg(&ctx->ac, ctx->args->oc_lds),
- 0, ac_glc, false);
+ 0, ac_glc);
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,
ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, inner_vec,
inner_comps, tf_inner_offset,
ac_get_arg(&ctx->ac, ctx->args->oc_lds),
- 0, ac_glc, false);
+ 0, ac_glc);
}
}
ac_optimize_vs_outputs(&ctx->ac,
ctx->main_function,
outinfo->vs_output_param_offset,
- VARYING_SLOT_MAX,
+ VARYING_SLOT_MAX, 0,
&outinfo->param_exports);
}
{
if (ctx->args->options->chip_class <= GFX8 &&
(ctx->stage == MESA_SHADER_GEOMETRY ||
- ctx->args->options->key.vs_common_out.as_es || ctx->args->options->key.vs_common_out.as_es)) {
+ ctx->args->options->key.vs_common_out.as_es)) {
unsigned ring = ctx->stage == MESA_SHADER_GEOMETRY ? RING_ESGS_GS
: RING_ESGS_VS;
LLVMValueRef offset = LLVMConstInt(ctx->ac.i32, ring, false);
LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm,
struct nir_shader *const *shaders,
int shader_count,
- struct radv_shader_info *shader_info,
- const struct radv_nir_compiler_options *options)
+ const struct radv_shader_args *args)
{
struct radv_shader_context ctx = {0};
- struct radv_shader_args args = {0};
- args.options = options;
- args.shader_info = shader_info;
- ctx.args = &args;
-
- declare_inputs(&args, shaders[shader_count - 1]->info.stage, shader_count >= 2,
- shader_count >= 2 ? shaders[shader_count - 2]->info.stage : MESA_SHADER_VERTEX);
+ ctx.args = args;
enum ac_float_mode float_mode = AC_FLOAT_MODE_DEFAULT;
- if (shader_info->float_controls_mode & FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP32) {
+ if (args->shader_info->float_controls_mode & FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP32) {
float_mode = AC_FLOAT_MODE_DENORM_FLUSH_TO_ZERO;
}
- ac_llvm_context_init(&ctx.ac, ac_llvm, options->chip_class,
- options->family, float_mode, shader_info->wave_size, 64);
+ ac_llvm_context_init(&ctx.ac, ac_llvm, args->options->chip_class,
+ args->options->family, float_mode,
+ args->shader_info->wave_size,
+ args->shader_info->ballot_bit_size);
ctx.context = ctx.ac.context;
ctx.max_workgroup_size = 0;
for (int i = 0; i < shader_count; ++i) {
ctx.max_workgroup_size = MAX2(ctx.max_workgroup_size,
- radv_nir_get_max_workgroup_size(args.options->chip_class,
+ radv_nir_get_max_workgroup_size(args->options->chip_class,
shaders[i]->info.stage,
shaders[i]));
}
if (ctx.ac.chip_class >= GFX10) {
if (is_pre_gs_stage(shaders[0]->info.stage) &&
- options->key.vs_common_out.as_ngg) {
+ args->options->key.vs_common_out.as_ngg) {
ctx.max_workgroup_size = 128;
}
}
ctx.abi.inputs = &ctx.inputs[0];
ctx.abi.emit_outputs = handle_shader_outputs_post;
- ctx.abi.emit_vertex = visit_emit_vertex;
+ ctx.abi.emit_vertex_with_counter = visit_emit_vertex_with_counter;
ctx.abi.load_ubo = radv_load_ubo;
ctx.abi.load_ssbo = radv_load_ssbo;
ctx.abi.load_sampler_desc = radv_get_sampler_desc;
ctx.abi.load_resource = radv_load_resource;
ctx.abi.clamp_shadow_reference = false;
- ctx.abi.robust_buffer_access = options->robust_buffer_access;
+ ctx.abi.robust_buffer_access = args->options->robust_buffer_access;
- bool is_ngg = is_pre_gs_stage(shaders[0]->info.stage) && args.options->key.vs_common_out.as_ngg;
+ bool is_ngg = is_pre_gs_stage(shaders[0]->info.stage) && args->options->key.vs_common_out.as_ngg;
if (shader_count >= 2 || is_ngg)
ac_init_exec_full_mask(&ctx.ac);
- if (args.ac.vertex_id.used)
- ctx.abi.vertex_id = ac_get_arg(&ctx.ac, args.ac.vertex_id);
- if (args.rel_auto_id.used)
- ctx.rel_auto_id = ac_get_arg(&ctx.ac, args.rel_auto_id);
- if (args.ac.instance_id.used)
- ctx.abi.instance_id = ac_get_arg(&ctx.ac, args.ac.instance_id);
+ if (args->ac.vertex_id.used)
+ ctx.abi.vertex_id = ac_get_arg(&ctx.ac, args->ac.vertex_id);
+ if (args->rel_auto_id.used)
+ ctx.rel_auto_id = ac_get_arg(&ctx.ac, args->rel_auto_id);
+ if (args->ac.instance_id.used)
+ ctx.abi.instance_id = ac_get_arg(&ctx.ac, args->ac.instance_id);
- if (options->has_ls_vgpr_init_bug &&
+ if (args->options->has_ls_vgpr_init_bug &&
shaders[shader_count - 1]->info.stage == MESA_SHADER_TESS_CTRL)
ac_nir_fixup_ls_hs_input_vgprs(&ctx);
* Add an extra dword per vertex to ensure an odd stride, which
* avoids bank conflicts for SoA accesses.
*/
- declare_esgs_ring(&ctx);
+ if (!args->options->key.vs_common_out.as_ngg_passthrough)
+ declare_esgs_ring(&ctx);
/* This is really only needed when streamout and / or vertex
* compaction is enabled.
*/
- LLVMTypeRef asi32 = LLVMArrayType(ctx.ac.i32, 8);
- ctx.gs_ngg_scratch = LLVMAddGlobalInAddressSpace(ctx.ac.module,
- asi32, "ngg_scratch", AC_ADDR_SPACE_LDS);
- LLVMSetInitializer(ctx.gs_ngg_scratch, LLVMGetUndef(asi32));
- LLVMSetAlignment(ctx.gs_ngg_scratch, 4);
+ if (args->shader_info->so.num_outputs) {
+ LLVMTypeRef asi32 = LLVMArrayType(ctx.ac.i32, 8);
+ ctx.gs_ngg_scratch = LLVMAddGlobalInAddressSpace(ctx.ac.module,
+ asi32, "ngg_scratch", AC_ADDR_SPACE_LDS);
+ LLVMSetInitializer(ctx.gs_ngg_scratch, LLVMGetUndef(asi32));
+ LLVMSetAlignment(ctx.gs_ngg_scratch, 4);
+ }
}
for(int i = 0; i < shader_count; ++i) {
ctx.gs_next_vertex[i] =
ac_build_alloca(&ctx.ac, ctx.ac.i32, "");
}
- if (args.options->key.vs_common_out.as_ngg) {
+ if (args->options->key.vs_common_out.as_ngg) {
for (unsigned i = 0; i < 4; ++i) {
ctx.gs_curprim_verts[i] =
ac_build_alloca(&ctx.ac, ctx.ac.i32, "");
}
unsigned scratch_size = 8;
- if (args.shader_info->so.num_outputs)
+ if (args->shader_info->so.num_outputs)
scratch_size = 44;
LLVMTypeRef ai32 = LLVMArrayType(ctx.ac.i32, scratch_size);
ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
ctx.abi.store_tcs_outputs = store_tcs_output;
if (shader_count == 1)
- ctx.tcs_num_inputs = args.options->key.tcs.num_inputs;
+ ctx.tcs_num_inputs = args->options->key.tcs.num_inputs;
else
- ctx.tcs_num_inputs = util_last_bit64(shader_info->vs.ls_outputs_written);
- ctx.tcs_num_patches = get_tcs_num_patches(&ctx);
+ ctx.tcs_num_inputs = util_last_bit64(args->shader_info->vs.ls_outputs_written);
+ unsigned tcs_num_outputs = util_last_bit64(ctx.args->shader_info->tcs.outputs_written);
+ unsigned tcs_num_patch_outputs = util_last_bit64(ctx.args->shader_info->tcs.patch_outputs_written);
+ ctx.tcs_num_patches =
+ get_tcs_num_patches(
+ ctx.args->options->key.tcs.input_vertices,
+ ctx.shader->info.tess.tcs_vertices_out,
+ ctx.tcs_num_inputs,
+ tcs_num_outputs,
+ tcs_num_patch_outputs,
+ ctx.args->options->tess_offchip_block_dw_size,
+ ctx.args->options->chip_class,
+ ctx.args->options->family);
} else if (shaders[i]->info.stage == MESA_SHADER_TESS_EVAL) {
ctx.abi.load_tess_varyings = load_tes_input;
ctx.abi.load_tess_coord = load_tess_coord;
ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
- ctx.tcs_num_patches = args.options->key.tes.num_patches;
+ ctx.tcs_num_patches = args->options->key.tes.num_patches;
} else if (shaders[i]->info.stage == MESA_SHADER_VERTEX) {
ctx.abi.load_base_vertex = radv_load_base_vertex;
} else if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT) {
ctx.abi.load_sample_position = load_sample_position;
ctx.abi.load_sample_mask_in = load_sample_mask_in;
- ctx.abi.emit_kill = radv_emit_kill;
}
if (shaders[i]->info.stage == MESA_SHADER_VERTEX &&
- args.options->key.vs_common_out.as_ngg &&
- args.options->key.vs_common_out.export_prim_id) {
+ args->options->key.vs_common_out.as_ngg &&
+ args->options->key.vs_common_out.export_prim_id) {
declare_esgs_ring(&ctx);
}
if (i) {
if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY &&
- args.options->key.vs_common_out.as_ngg) {
+ args->options->key.vs_common_out.as_ngg) {
gfx10_ngg_gs_emit_prologue(&ctx);
nested_barrier = false;
} else {
ac_setup_rings(&ctx);
- LLVMBasicBlockRef merge_block;
+ LLVMBasicBlockRef merge_block = NULL;
if (shader_count >= 2 || is_ngg) {
LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx.ac.builder));
LLVMBasicBlockRef then_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
LLVMValueRef count =
ac_unpack_param(&ctx.ac,
- ac_get_arg(&ctx.ac, args.merged_wave_info),
+ ac_get_arg(&ctx.ac, args->merged_wave_info),
8 * i, 8);
LLVMValueRef thread_id = ac_get_thread_id(&ctx.ac);
LLVMValueRef cond = LLVMBuildICmp(ctx.ac.builder, LLVMIntULT,
else if(shaders[i]->info.stage == MESA_SHADER_GEOMETRY)
prepare_gs_input_vgprs(&ctx, shader_count >= 2);
- ac_nir_translate(&ctx.ac, &ctx.abi, &args.ac, shaders[i]);
+ ac_nir_translate(&ctx.ac, &ctx.abi, &args->ac, shaders[i]);
if (shader_count >= 2 || is_ngg) {
LLVMBuildBr(ctx.ac.builder, merge_block);
/* This needs to be outside the if wrapping the shader body, as sometimes
* the HW generates waves with 0 es/vs threads. */
if (is_pre_gs_stage(shaders[i]->info.stage) &&
- args.options->key.vs_common_out.as_ngg &&
+ args->options->key.vs_common_out.as_ngg &&
i == shader_count - 1) {
handle_ngg_outputs_post_2(&ctx);
} else if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY &&
- args.options->key.vs_common_out.as_ngg) {
+ args->options->key.vs_common_out.as_ngg) {
gfx10_ngg_gs_emit_epilogue_2(&ctx);
}
if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
- shader_info->tcs.num_patches = ctx.tcs_num_patches;
- shader_info->tcs.lds_size = calculate_tess_lds_size(&ctx);
+ unsigned tcs_num_outputs = util_last_bit64(ctx.args->shader_info->tcs.outputs_written);
+ unsigned tcs_num_patch_outputs = util_last_bit64(ctx.args->shader_info->tcs.patch_outputs_written);
+ args->shader_info->tcs.num_patches = ctx.tcs_num_patches;
+ args->shader_info->tcs.num_lds_blocks =
+ calculate_tess_lds_size(
+ ctx.args->options->chip_class,
+ ctx.args->options->key.tcs.input_vertices,
+ ctx.shader->info.tess.tcs_vertices_out,
+ ctx.tcs_num_inputs,
+ ctx.tcs_num_patches,
+ tcs_num_outputs,
+ tcs_num_patch_outputs);
}
}
LLVMBuildRetVoid(ctx.ac.builder);
- if (options->dump_preoptir) {
+ if (args->options->dump_preoptir) {
fprintf(stderr, "%s LLVM IR:\n\n",
- radv_get_shader_name(shader_info,
+ radv_get_shader_name(args->shader_info,
shaders[shader_count - 1]->info.stage));
ac_dump_module(ctx.ac.module);
fprintf(stderr, "\n");
}
- ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
+ ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, args->options);
if (shader_count == 1)
ac_nir_eliminate_const_vs_outputs(&ctx);
- if (options->dump_shader) {
- args.shader_info->private_mem_vgprs =
+ if (args->options->dump_shader) {
+ args->shader_info->private_mem_vgprs =
ac_count_scratch_private_memory(ctx.main_function);
}
free(elf_buffer);
}
-void
+static void
radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm,
struct radv_shader_binary **rbinary,
- struct radv_shader_info *shader_info,
+ const struct radv_shader_args *args,
struct nir_shader *const *nir,
- int nir_count,
- const struct radv_nir_compiler_options *options)
+ int nir_count)
{
LLVMModuleRef llvm_module;
- llvm_module = ac_translate_nir_to_llvm(ac_llvm, nir, nir_count, shader_info,
- options);
+ llvm_module = ac_translate_nir_to_llvm(ac_llvm, nir, nir_count, args);
ac_compile_llvm_module(ac_llvm, llvm_module, rbinary,
nir[nir_count - 1]->info.stage,
- radv_get_shader_name(shader_info,
+ radv_get_shader_name(args->shader_info,
nir[nir_count - 1]->info.stage),
- options);
+ args->options);
/* Determine the ES type (VS or TES) for the GS on GFX9. */
- if (options->chip_class >= GFX9) {
+ if (args->options->chip_class >= GFX9) {
if (nir_count == 2 &&
nir[1]->info.stage == MESA_SHADER_GEOMETRY) {
- shader_info->gs.es_type = nir[0]->info.stage;
+ args->shader_info->gs.es_type = nir[0]->info.stage;
}
}
}
LLVMPositionBuilderAtEnd(ctx->ac.builder, end_bb);
}
-void
+static void
radv_compile_gs_copy_shader(struct ac_llvm_compiler *ac_llvm,
struct nir_shader *geom_shader,
struct radv_shader_binary **rbinary,
- struct radv_shader_info *shader_info,
- const struct radv_nir_compiler_options *options)
+ const struct radv_shader_args *args)
{
struct radv_shader_context ctx = {0};
- struct radv_shader_args args = {0};
- args.options = options;
- args.shader_info = shader_info;
- ctx.args = &args;
+ ctx.args = args;
- args.is_gs_copy_shader = true;
- declare_inputs(&args, MESA_SHADER_VERTEX, false, MESA_SHADER_VERTEX);
+ assert(args->is_gs_copy_shader);
- ac_llvm_context_init(&ctx.ac, ac_llvm, options->chip_class,
- options->family, AC_FLOAT_MODE_DEFAULT, 64, 64);
+ ac_llvm_context_init(&ctx.ac, ac_llvm, args->options->chip_class,
+ args->options->family, AC_FLOAT_MODE_DEFAULT, 64, 64);
ctx.context = ctx.ac.context;
ctx.stage = MESA_SHADER_VERTEX;
LLVMBuildRetVoid(ctx.ac.builder);
- ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
+ ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, args->options);
ac_compile_llvm_module(ac_llvm, ctx.ac.module, rbinary,
- MESA_SHADER_VERTEX, "GS Copy Shader", options);
+ MESA_SHADER_VERTEX, "GS Copy Shader", args->options);
(*rbinary)->is_gs_copy_shader = true;
}
+
+void
+llvm_compile_shader(struct radv_device *device,
+ unsigned shader_count,
+ struct nir_shader *const *shaders,
+ struct radv_shader_binary **binary,
+ struct radv_shader_args *args)
+{
+ enum ac_target_machine_options tm_options = 0;
+ struct ac_llvm_compiler ac_llvm;
+ bool thread_compiler;
+
+ tm_options |= AC_TM_SUPPORTS_SPILL;
+ if (args->options->check_ir)
+ tm_options |= AC_TM_CHECK_IR;
+
+ thread_compiler = !(device->instance->debug_flags & RADV_DEBUG_NOTHREADLLVM);
+
+ radv_init_llvm_compiler(&ac_llvm, thread_compiler,
+ args->options->family, tm_options,
+ args->shader_info->wave_size);
+
+ if (args->is_gs_copy_shader) {
+ radv_compile_gs_copy_shader(&ac_llvm, *shaders, binary, args);
+ } else {
+ radv_compile_nir_shader(&ac_llvm, binary, args,
+ shaders, shader_count);
+ }
+
+ radv_destroy_llvm_compiler(&ac_llvm, thread_compiler);
+}
projects / mesa.git / blobdiff