AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

ac/nir/ngg,radv: move LDS layout calculation out of nir ngg lowering

Browse Source 
Use lds base load intrinsics in nir ngg lowering to get layout, left
its calulation to driver.

Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Signed-off-by: Qiang Yu <yuq825@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18832>
This commit is contained in:
Qiang Yu
2022-07-22 17:42:41 +08:00
committed by Marge Bot
parent 3d6cce2e4c
commit e536d0fe4b
6 changed files with 226 additions and 84 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/amd/common/ac_nir.h
-3
View File
@@ -122,7 +122,6 @@ ac_nir_lower_indirect_derefs(nir_shader *shader,
void
ac_nir_lower_ngg_nogs(nir_shader *shader,
enum radeon_family family,
unsigned max_num_es_vertices,
unsigned num_vertices_per_primitive,
unsigned max_workgroup_size,
unsigned wave_size,
@@ -142,9 +141,7 @@ ac_nir_lower_ngg_gs(nir_shader *shader,
enum amd_gfx_level gfx_level,
unsigned wave_size,
unsigned max_workgroup_size,
unsigned esgs_ring_lds_bytes,
unsigned gs_out_vtx_bytes,
unsigned gs_total_out_vtx_bytes,
bool has_xfb_query,
bool can_cull,
bool disable_streamout);
src/amd/common/ac_nir_lower_ngg.c
+147 -77
View File
@@ -66,13 +66,11 @@ typedef struct
unsigned wave_size;
unsigned max_num_waves;
unsigned num_vertices_per_primitives;
unsigned max_es_num_vertices;
unsigned position_store_base;
int primitive_id_location;
/* LDS params */
unsigned pervertex_lds_bytes;
unsigned total_lds_bytes;
uint64_t inputs_needed_by_pos;
uint64_t inputs_needed_by_others;
@@ -110,8 +108,8 @@ typedef struct
unsigned wave_size;
unsigned max_num_waves;
unsigned num_vertices_per_primitive;
unsigned lds_addr_gs_out_vtx;
unsigned lds_addr_gs_scratch;
nir_ssa_def *lds_addr_gs_out_vtx;
nir_ssa_def *lds_addr_gs_scratch;
unsigned lds_bytes_per_gs_out_vertex;
unsigned lds_offs_primflags;
bool has_xfb_query;
@@ -303,7 +301,7 @@ summarize_repack(nir_builder *b, nir_ssa_def *packed_counts, unsigned num_lds_dw
*/
static wg_repack_result
repack_invocations_in_workgroup(nir_builder *b, nir_ssa_def *input_bool,
unsigned lds_addr_base, unsigned max_num_waves,
nir_ssa_def *lds_addr_base, unsigned max_num_waves,
unsigned wave_size)
{
/* Input boolean: 1 if the current invocation should survive the repack. */
@@ -340,15 +338,17 @@ repack_invocations_in_workgroup(nir_builder *b, nir_ssa_def *input_bool,
assert(num_lds_dwords <= 2);
nir_ssa_def *wave_id = nir_load_subgroup_id(b);
nir_ssa_def *lds_offset = nir_iadd(b, lds_addr_base, wave_id);
nir_ssa_def *dont_care = nir_ssa_undef(b, 1, num_lds_dwords * 32);
nir_if *if_first_lane = nir_push_if(b, nir_elect(b, 1));
nir_store_shared(b, nir_u2u8(b, surviving_invocations_in_current_wave), wave_id, .base = lds_addr_base);
nir_store_shared(b, nir_u2u8(b, surviving_invocations_in_current_wave), lds_offset);
nir_scoped_barrier(b, .execution_scope=NIR_SCOPE_WORKGROUP, .memory_scope=NIR_SCOPE_WORKGROUP,
.memory_semantics=NIR_MEMORY_ACQ_REL, .memory_modes=nir_var_mem_shared);
nir_ssa_def *packed_counts = nir_load_shared(b, 1, num_lds_dwords * 32, nir_imm_int(b, 0), .base = lds_addr_base, .align_mul = 8u);
nir_ssa_def *packed_counts =
nir_load_shared(b, 1, num_lds_dwords * 32, lds_addr_base, .align_mul = 8u);
nir_pop_if(b, if_first_lane);
@@ -893,7 +893,6 @@ compact_vertices_after_culling(nir_builder *b,
nir_ssa_def *es_exporter_tid,
nir_ssa_def *num_live_vertices_in_workgroup,
nir_ssa_def *fully_culled,
unsigned ngg_scratch_lds_base_addr,
unsigned pervertex_lds_bytes,
unsigned max_exported_args)
{
@@ -1273,25 +1272,38 @@ clipdist_culling_es_part(nir_builder *b, lower_ngg_nogs_state *nogs_state,
}
}
static unsigned
ngg_nogs_get_culling_pervertex_lds_size(gl_shader_stage stage,
bool uses_instance_id,
bool uses_primitive_id,
unsigned *max_exported_args)
{
unsigned max_args;
if (stage == MESA_SHADER_VERTEX) {
max_args = uses_instance_id ? 2 : 1;
} else {
assert(stage == MESA_SHADER_TESS_EVAL);
max_args = uses_primitive_id ? 4 : 3;
}
if (max_exported_args)
*max_exported_args = max_args;
return lds_es_arg_0 + max_args * 4u;
}
static void
add_deferred_attribute_culling(nir_builder *b, nir_cf_list *original_extracted_cf, lower_ngg_nogs_state *nogs_state)
{
bool uses_instance_id = BITSET_TEST(b->shader->info.system_values_read, SYSTEM_VALUE_INSTANCE_ID);
bool uses_tess_primitive_id = BITSET_TEST(b->shader->info.system_values_read, SYSTEM_VALUE_PRIMITIVE_ID);
unsigned max_exported_args = b->shader->info.stage == MESA_SHADER_VERTEX ? 2 : 4;
if (b->shader->info.stage == MESA_SHADER_VERTEX && !uses_instance_id)
max_exported_args--;
else if (b->shader->info.stage == MESA_SHADER_TESS_EVAL && !uses_tess_primitive_id)
max_exported_args--;
unsigned pervertex_lds_bytes = lds_es_arg_0 + max_exported_args * 4u;
unsigned total_es_lds_bytes = pervertex_lds_bytes * nogs_state->max_es_num_vertices;
unsigned max_num_waves = nogs_state->max_num_waves;
unsigned ngg_scratch_lds_base_addr = ALIGN(total_es_lds_bytes, 8u);
unsigned ngg_scratch_lds_bytes = ALIGN(max_num_waves, 4u);
nogs_state->total_lds_bytes = MAX2(nogs_state->total_lds_bytes,
ngg_scratch_lds_base_addr + ngg_scratch_lds_bytes);
unsigned max_exported_args;
unsigned pervertex_lds_bytes =
ngg_nogs_get_culling_pervertex_lds_size(b->shader->info.stage,
uses_instance_id,
uses_tess_primitive_id,
&max_exported_args);
nir_function_impl *impl = nir_shader_get_entrypoint(b->shader);
@@ -1368,6 +1380,8 @@ add_deferred_attribute_culling(nir_builder *b, nir_cf_list *original_extracted_c
remove_culling_shader_outputs(b->shader, nogs_state);
b->cursor = nir_after_cf_list(&impl->body);
nir_ssa_def *lds_scratch_base = nir_load_lds_ngg_scratch_base_amd(b);
/* Run culling algorithms if culling is enabled.
*
* NGG culling can be enabled or disabled in runtime.
@@ -1470,8 +1484,9 @@ add_deferred_attribute_culling(nir_builder *b, nir_cf_list *original_extracted_c
nir_ssa_def *es_accepted = nir_load_var(b, es_accepted_var);
/* Repack the vertices that survived the culling. */
wg_repack_result rep = repack_invocations_in_workgroup(b, es_accepted, ngg_scratch_lds_base_addr,
nogs_state->max_num_waves, nogs_state->wave_size);
wg_repack_result rep = repack_invocations_in_workgroup(b, es_accepted, lds_scratch_base,
nogs_state->max_num_waves,
nogs_state->wave_size);
nir_ssa_def *num_live_vertices_in_workgroup = rep.num_repacked_invocations;
nir_ssa_def *es_exporter_tid = rep.repacked_invocation_index;
@@ -1492,7 +1507,7 @@ add_deferred_attribute_culling(nir_builder *b, nir_cf_list *original_extracted_c
repacked_arg_vars, gs_vtxaddr_vars,
invocation_index, es_vertex_lds_addr,
es_exporter_tid, num_live_vertices_in_workgroup, fully_culled,
ngg_scratch_lds_base_addr, pervertex_lds_bytes, max_exported_args);
pervertex_lds_bytes, max_exported_args);
}
nir_push_else(b, if_cull_en);
{
@@ -1601,8 +1616,8 @@ ngg_nogs_store_all_outputs_to_lds(nir_shader *shader, lower_ngg_nogs_state *st)
static void
ngg_build_streamout_buffer_info(nir_builder *b,
nir_xfb_info *info,
unsigned scratch_base,
bool has_prim_query,
nir_ssa_def *scratch_base,
nir_ssa_def *tid_in_tg,
nir_ssa_def *gen_prim[4],
nir_ssa_def *prim_stride_ret[4],
@@ -1677,8 +1692,7 @@ ngg_build_streamout_buffer_info(nir_builder *b,
nir_imin(b, emit_prim[stream], remain_prim));
/* Save to LDS for being accessed by other waves in this workgroup. */
nir_store_shared(b, buffer_offset, nir_imm_int(b, buffer * 4),
.base = scratch_base);
nir_store_shared(b, buffer_offset, scratch_base, .base = buffer * 4);
}
/* No need to fixup the global buffer offset once we overflowed,
@@ -1690,8 +1704,7 @@ ngg_build_streamout_buffer_info(nir_builder *b,
if (!(info->streams_written & BITFIELD_BIT(stream)))
continue;
nir_store_shared(b, emit_prim[stream], nir_imm_int(b, stream * 4),
.base = scratch_base + 16);
nir_store_shared(b, emit_prim[stream], scratch_base, .base = 16 + stream * 4);
}
/* Update shader query. */
@@ -1719,7 +1732,7 @@ ngg_build_streamout_buffer_info(nir_builder *b,
continue;
buffer_offsets_ret[buffer] =
nir_load_shared(b, 1, 32, nir_imm_int(b, buffer * 4), .base = scratch_base);
nir_load_shared(b, 1, 32, scratch_base, .base = buffer * 4);
}
/* Fetch the per-stream emit prim in all waves. */
@@ -1728,7 +1741,7 @@ ngg_build_streamout_buffer_info(nir_builder *b,
continue;
emit_prim_ret[stream] =
nir_load_shared(b, 1, 32, nir_imm_int(b, stream * 4), .base = scratch_base + 16);
nir_load_shared(b, 1, 32, scratch_base, .base = 16 + stream * 4);
}
}
@@ -1780,11 +1793,7 @@ ngg_nogs_build_streamout(nir_builder *b, lower_ngg_nogs_state *s)
return;
}
unsigned total_es_lds_bytes = s->pervertex_lds_bytes * s->max_es_num_vertices;
unsigned scratch_base = ALIGN(total_es_lds_bytes, 8u);
/* 4 dwords for 4 streamout buffer offset, 1 dword for emit prim count */
unsigned scratch_size = 20;
s->total_lds_bytes = MAX2(s->total_lds_bytes, scratch_base + scratch_size);
nir_ssa_def *lds_scratch_base = nir_load_lds_ngg_scratch_base_amd(b);
/* Get global buffer offset where this workgroup will stream out data to. */
nir_ssa_def *generated_prim = nir_load_workgroup_num_input_primitives_amd(b);
@@ -1794,8 +1803,8 @@ ngg_nogs_build_streamout(nir_builder *b, lower_ngg_nogs_state *s)
nir_ssa_def *so_buffer[4] = {0};
nir_ssa_def *prim_stride[4] = {0};
nir_ssa_def *tid_in_tg = nir_load_local_invocation_index(b);
ngg_build_streamout_buffer_info(b, info, scratch_base,
s->has_prim_query, tid_in_tg,
ngg_build_streamout_buffer_info(b, info, s->has_prim_query,
lds_scratch_base, tid_in_tg,
gen_prim_per_stream, prim_stride,
so_buffer, buffer_offsets,
emit_prim_per_stream);
@@ -1840,10 +1849,42 @@ ngg_nogs_build_streamout(nir_builder *b, lower_ngg_nogs_state *s)
free(info);
}
static unsigned
ngg_nogs_get_pervertex_lds_size(gl_shader_stage stage,
unsigned shader_num_outputs,
bool streamout_enabled,
bool export_prim_id,
bool has_user_edgeflags)
{
unsigned pervertex_lds_bytes = 0;
if (streamout_enabled) {
/* The extra dword is used to avoid LDS bank conflicts and store the primitive id.
* TODO: only alloc space for outputs that really need streamout.
*/
pervertex_lds_bytes = (shader_num_outputs * 4 + 1) * 4;
}
bool need_prim_id_store_shared = export_prim_id && stage == MESA_SHADER_VERTEX;
if (need_prim_id_store_shared || has_user_edgeflags) {
unsigned size = 0;
if (need_prim_id_store_shared)
size += 4;
if (has_user_edgeflags)
size += 4;
/* pad to odd dwords to avoid LDS bank conflict */
size |= 4;
pervertex_lds_bytes = MAX2(pervertex_lds_bytes, size);
}
return pervertex_lds_bytes;
}
void
ac_nir_lower_ngg_nogs(nir_shader *shader,
enum radeon_family family,
unsigned max_num_es_vertices,
unsigned num_vertices_per_primitives,
unsigned max_workgroup_size,
unsigned wave_size,
@@ -1860,7 +1901,7 @@ ac_nir_lower_ngg_nogs(nir_shader *shader,
{
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
assert(impl);
assert(max_num_es_vertices && max_workgroup_size && wave_size);
assert(max_workgroup_size && wave_size);
assert(!(can_cull && passthrough));
nir_variable *position_value_var = nir_local_variable_create(impl, glsl_vec4_type(), "position_value");
@@ -1894,7 +1935,6 @@ ac_nir_lower_ngg_nogs(nir_shader *shader,
.es_accepted_var = es_accepted_var,
.gs_accepted_var = gs_accepted_var,
.max_num_waves = DIV_ROUND_UP(max_workgroup_size, wave_size),
.max_es_num_vertices = max_num_es_vertices,
.wave_size = wave_size,
.instance_rate_inputs = instance_rate_inputs,
.clipdist_enable_mask = clipdist_enable_mask,
@@ -1965,23 +2005,12 @@ ac_nir_lower_ngg_nogs(nir_shader *shader,
}
/* determine the LDS vertex stride */
if (state.streamout_enabled) {
/* The extra dword is used to avoid LDS bank conflicts and store the primitive id.
* TODO: only alloc space for outputs that really need streamout.
*/
state.pervertex_lds_bytes = (shader->num_outputs * 4 + 1) * 4;
} else if (need_prim_id_store_shared || state.has_user_edgeflags) {
if (need_prim_id_store_shared)
state.pervertex_lds_bytes += 4;
if (state.has_user_edgeflags)
state.pervertex_lds_bytes += 4;
/* pad to odd dwords to avoid LDS bank conflict */
state.pervertex_lds_bytes |= 4;
state.total_lds_bytes = MAX2(state.total_lds_bytes,
state.pervertex_lds_bytes * max_num_es_vertices);
}
state.pervertex_lds_bytes =
ngg_nogs_get_pervertex_lds_size(shader->info.stage,
shader->num_outputs,
state.streamout_enabled,
primitive_id_location >= 0,
state.has_user_edgeflags);
if (need_prim_id_store_shared) {
emit_ngg_nogs_prim_id_store_shared(b, &state);
@@ -2071,8 +2100,6 @@ ac_nir_lower_ngg_nogs(nir_shader *shader,
if (can_cull)
progress |= cleanup_culling_shader_after_dce(shader, b->impl, &state);
} while (progress);
shader->info.shared_size = state.total_lds_bytes;
}
/**
@@ -2119,7 +2146,7 @@ ngg_gs_out_vertex_addr(nir_builder *b, nir_ssa_def *out_vtx_idx, lower_ngg_gs_st
}
nir_ssa_def *out_vtx_offs = nir_imul_imm(b, out_vtx_idx, s->lds_bytes_per_gs_out_vertex);
return nir_iadd_imm_nuw(b, out_vtx_offs, s->lds_addr_gs_out_vtx);
return nir_iadd_nuw(b, out_vtx_offs, s->lds_addr_gs_out_vtx);
}
static nir_ssa_def *
@@ -2729,6 +2756,8 @@ ngg_gs_build_streamout(nir_builder *b, lower_ngg_gs_state *st)
prim_live[stream] = nir_i2b(b, nir_iand_imm(b, out_vtx_primflag[stream], 1));
unsigned scratch_stride = ALIGN(st->max_num_waves, 4);
nir_ssa_def *scratch_base =
nir_iadd_imm(b, st->lds_addr_gs_scratch, stream * scratch_stride);
/* We want to export primitives to streamout buffer in sequence,
* but not all vertices are alive or mark end of a primitive, so
@@ -2743,8 +2772,7 @@ ngg_gs_build_streamout(nir_builder *b, lower_ngg_gs_state *st)
* stream..
*/
wg_repack_result rep =
repack_invocations_in_workgroup(b, prim_live[stream],
st->lds_addr_gs_scratch + stream * scratch_stride,
repack_invocations_in_workgroup(b, prim_live[stream], scratch_base,
st->max_num_waves, st->wave_size);
/* nir_intrinsic_set_vertex_and_primitive_count can also get primitive count of
@@ -2766,8 +2794,8 @@ ngg_gs_build_streamout(nir_builder *b, lower_ngg_gs_state *st)
nir_ssa_def *buffer_offsets[4] = {0};
nir_ssa_def *so_buffer[4] = {0};
nir_ssa_def *prim_stride[4] = {0};
ngg_build_streamout_buffer_info(b, info, st->lds_addr_gs_scratch,
st->has_xfb_query, tid_in_tg, gen_prim,
ngg_build_streamout_buffer_info(b, info, st->has_xfb_query,
st->lds_addr_gs_scratch, tid_in_tg, gen_prim,
prim_stride, so_buffer, buffer_offsets, emit_prim);
/* GS use packed location for vertex LDS storage. */
@@ -2884,9 +2912,7 @@ ac_nir_lower_ngg_gs(nir_shader *shader,
enum amd_gfx_level gfx_level,
unsigned wave_size,
unsigned max_workgroup_size,
unsigned esgs_ring_lds_bytes,
unsigned gs_out_vtx_bytes,
unsigned gs_total_out_vtx_bytes,
bool has_xfb_query,
bool can_cull,
bool disable_streamout)
@@ -2899,8 +2925,6 @@ ac_nir_lower_ngg_gs(nir_shader *shader,
.gfx_level = gfx_level,
.max_num_waves = DIV_ROUND_UP(max_workgroup_size, wave_size),
.wave_size = wave_size,
.lds_addr_gs_out_vtx = esgs_ring_lds_bytes,
.lds_addr_gs_scratch = ALIGN(esgs_ring_lds_bytes + gs_total_out_vtx_bytes, 8u /* for the repacking code */),
.lds_offs_primflags = gs_out_vtx_bytes,
.lds_bytes_per_gs_out_vertex = gs_out_vtx_bytes + 4u,
.can_cull = can_cull,
@@ -2908,14 +2932,6 @@ ac_nir_lower_ngg_gs(nir_shader *shader,
.has_xfb_query = has_xfb_query,
};
unsigned lds_scratch_bytes = ALIGN(state.max_num_waves, 4u);
/* streamout take 8 dwords for buffer offset and emit vertex per stream */
if (state.streamout_enabled)
lds_scratch_bytes = MAX2(lds_scratch_bytes, 32);
unsigned total_lds_bytes = state.lds_addr_gs_scratch + lds_scratch_bytes;
shader->info.shared_size = total_lds_bytes;
if (!can_cull) {
nir_gs_count_vertices_and_primitives(shader, state.const_out_vtxcnt,
state.const_out_prmcnt, 4u);
@@ -2949,6 +2965,9 @@ ac_nir_lower_ngg_gs(nir_shader *shader,
nir_scoped_barrier(b, .execution_scope=NIR_SCOPE_WORKGROUP, .memory_scope=NIR_SCOPE_WORKGROUP,
.memory_semantics=NIR_MEMORY_ACQ_REL, .memory_modes=nir_var_mem_shared);
state.lds_addr_gs_out_vtx = nir_load_lds_ngg_gs_out_vertex_base_amd(b);
state.lds_addr_gs_scratch = nir_load_lds_ngg_scratch_base_amd(b);
/* Wrap the GS control flow. */
nir_if *if_gs_thread = nir_push_if(b, nir_has_input_primitive_amd(b));
@@ -2982,6 +3001,57 @@ ac_nir_lower_ngg_gs(nir_shader *shader,
nir_metadata_preserve(impl, nir_metadata_none);
}
unsigned
ac_ngg_nogs_get_pervertex_lds_size(gl_shader_stage stage,
unsigned shader_num_outputs,
bool streamout_enabled,
bool export_prim_id,
bool has_user_edgeflags,
bool can_cull,
bool uses_instance_id,
bool uses_primitive_id)
{
/* for culling time lds layout only */
unsigned culling_pervertex_lds_bytes = can_cull ?
ngg_nogs_get_culling_pervertex_lds_size(
stage, uses_instance_id, uses_primitive_id, NULL) : 0;
unsigned pervertex_lds_bytes =
ngg_nogs_get_pervertex_lds_size(stage, shader_num_outputs, streamout_enabled,
export_prim_id, has_user_edgeflags);
return MAX2(culling_pervertex_lds_bytes, pervertex_lds_bytes);
}
unsigned
ac_ngg_get_scratch_lds_size(gl_shader_stage stage,
unsigned workgroup_size,
unsigned wave_size,
bool streamout_enabled,
bool can_cull)
{
unsigned scratch_lds_size = 0;
unsigned max_num_waves = DIV_ROUND_UP(workgroup_size, wave_size);
if (stage == MESA_SHADER_VERTEX || stage == MESA_SHADER_TESS_EVAL) {
if (streamout_enabled) {
/* 4 dwords for 4 streamout buffer offset, 1 dword for emit prim count */
scratch_lds_size = 20;
} else if (can_cull) {
scratch_lds_size = ALIGN(max_num_waves, 4u);
}
} else {
assert(stage == MESA_SHADER_GEOMETRY);
scratch_lds_size = ALIGN(max_num_waves, 4u);
/* streamout take 8 dwords for buffer offset and emit vertex per stream */
if (streamout_enabled)
scratch_lds_size = MAX2(scratch_lds_size, 32);
}
return scratch_lds_size;
}
static void
ms_store_prim_indices(nir_builder *b,
nir_ssa_def *val,
src/amd/common/ac_shader_util.h
+17
View File
@@ -170,6 +170,23 @@ void ac_get_scratch_tmpring_size(const struct radeon_info *info, bool compute,
unsigned bytes_per_wave, unsigned *max_seen_bytes_per_wave,
uint32_t *tmpring_size);
unsigned
ac_ngg_nogs_get_pervertex_lds_size(gl_shader_stage stage,
unsigned shader_num_outputs,
bool streamout_enabled,
bool export_prim_id,
bool has_user_edgeflags,
bool can_cull,
bool uses_instance_id,
bool uses_primitive_id);
unsigned
ac_ngg_get_scratch_lds_size(gl_shader_stage stage,
unsigned workgroup_size,
unsigned wave_size,
bool streamout_enabled,
bool can_cull);
#ifdef __cplusplus
}
#endif
src/amd/vulkan/radv_nir_lower_abi.c
+7
View File
@@ -370,6 +370,13 @@ lower_abi_instr(nir_builder *b, nir_instr *instr, void *state)
replacement =
ac_nir_load_arg(b, &s->args->ac, s->args->ac.streamout_offset[nir_intrinsic_base(intrin)]);
break;
case nir_intrinsic_load_lds_ngg_gs_out_vertex_base_amd:
replacement = nir_imm_int(b, s->info->ngg_info.esgs_ring_size);
break;
case nir_intrinsic_load_lds_ngg_scratch_base_amd:
replacement = nir_imm_int(b, s->info->ngg_info.scratch_lds_base);
break;
default:
progress = false;
break;
src/amd/vulkan/radv_shader.c
+54 -4
View File
@@ -1326,6 +1326,56 @@ radv_consider_culling(const struct radv_physical_device *pdevice, struct nir_sha
return true;
}
static void
setup_ngg_lds_layout(nir_shader *nir, struct radv_shader_info *info,
unsigned max_vtx_in)
{
unsigned scratch_lds_base = 0;
gl_shader_stage stage = nir->info.stage;
if (stage == MESA_SHADER_VERTEX || stage == MESA_SHADER_TESS_EVAL) {
/* Get pervertex LDS usage. */
bool uses_instanceid =
BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_INSTANCE_ID);
bool uses_primtive_id =
BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_PRIMITIVE_ID);
unsigned pervertex_lds_bytes =
ac_ngg_nogs_get_pervertex_lds_size(stage,
nir->num_outputs,
false, /* streamout */
info->outinfo.export_prim_id,
false, /* user edge flag */
info->has_ngg_culling,
uses_instanceid,
uses_primtive_id);
unsigned total_es_lds_bytes = pervertex_lds_bytes * max_vtx_in;
scratch_lds_base = ALIGN(total_es_lds_bytes, 8u);
} else if (stage == MESA_SHADER_GEOMETRY) {
unsigned esgs_ring_lds_bytes = info->ngg_info.esgs_ring_size;
unsigned gs_total_out_vtx_bytes = info->ngg_info.ngg_emit_size * 4u;
scratch_lds_base =
ALIGN(esgs_ring_lds_bytes + gs_total_out_vtx_bytes, 8u /* for the repacking code */);
} else {
/* not handled here */
return;
}
/* Get scratch LDS usage. */
unsigned scratch_lds_size =
ac_ngg_get_scratch_lds_size(stage,
info->workgroup_size,
info->wave_size,
false, /* streamout */
info->has_ngg_culling);
/* Get total LDS usage. */
nir->info.shared_size = scratch_lds_base + scratch_lds_size;
/* Record scratch base for abi lower of nir_load_lds_ngg_scratch_base_amd. */
info->ngg_info.scratch_lds_base = scratch_lds_base;
}
void radv_lower_ngg(struct radv_device *device, struct radv_pipeline_stage *ngg_stage,
const struct radv_pipeline_key *pl_key)
{
@@ -1374,6 +1424,8 @@ void radv_lower_ngg(struct radv_device *device, struct radv_pipeline_stage *ngg_
/* Invocations that process an input vertex */
unsigned max_vtx_in = MIN2(256, ngg_info->hw_max_esverts);
setup_ngg_lds_layout(nir, &ngg_stage->info, max_vtx_in);
if (nir->info.stage == MESA_SHADER_VERTEX ||
nir->info.stage == MESA_SHADER_TESS_EVAL) {
bool export_prim_id = info->outinfo.export_prim_id;
@@ -1384,8 +1436,7 @@ void radv_lower_ngg(struct radv_device *device, struct radv_pipeline_stage *ngg_
radv_optimize_nir_algebraic(nir, false);
NIR_PASS_V(nir, ac_nir_lower_ngg_nogs,
device->physical_device->rad_info.family,
max_vtx_in, num_vertices_per_prim,
device->physical_device->rad_info.family, num_vertices_per_prim,
info->workgroup_size, info->wave_size, info->has_ngg_culling,
info->has_ngg_early_prim_export, info->is_ngg_passthrough,
false, pl_key->primitives_generated_query,
@@ -1399,8 +1450,7 @@ void radv_lower_ngg(struct radv_device *device, struct radv_pipeline_stage *ngg_
NIR_PASS_V(nir, ac_nir_lower_ngg_gs,
device->physical_device->rad_info.gfx_level,
info->wave_size, info->workgroup_size,
info->ngg_info.esgs_ring_size, info->gs.gsvs_vertex_size,
info->ngg_info.ngg_emit_size * 4u, true, false, true);
info->gs.gsvs_vertex_size, true, false, true);
} else if (nir->info.stage == MESA_SHADER_MESH) {
bool scratch_ring = false;
NIR_PASS_V(nir, ac_nir_lower_ngg_ms, &scratch_ring, info->wave_size, pl_key->has_multiview_view_index);
src/amd/vulkan/radv_shader.h
+1
View File
@@ -227,6 +227,7 @@ struct gfx10_ngg_info {
uint32_t prim_amp_factor;
uint32_t vgt_esgs_ring_itemsize;
uint32_t esgs_ring_size;
uint32_t scratch_lds_base;
bool max_vert_out_per_gs_instance;
};