From 65c5ee1628a62adad38b919f2f3d8bb3eec725e6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= Date: Wed, 11 Jun 2025 14:39:00 -0400 Subject: [PATCH] radeonsi: stop using LLVM LDS linking logic for the GS out LDS offset MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit This will enable large code removal. shader->config.lds_size is now always computed the same as ACO except for compute shaders. We have to add a new 8-bit user SGPR bitfield called GS_STATE_GS_OUT_LDS_OFFSET_256B, which contains the offset that was previously set by the relocation. Since the offset must be a multiple of 256, we have to add padding to the LDS size computation to make sure the alignment to 256 for the ESGS LDS size doesn't cause us to exceed the maximum LDS size. Acked-by: Pierre-Eric Pelloux-Prayer Reviewed-by: Timur Kristóf Part-of: --- src/amd/common/ac_shader_util.c | 4 +-- src/amd/common/ac_shader_util.h | 2 +- src/amd/vulkan/radv_shader_info.c | 2 +- .../drivers/radeonsi/si_nir_lower_abi.c | 3 +++ src/gallium/drivers/radeonsi/si_shader.c | 27 +++++++++++++------ src/gallium/drivers/radeonsi/si_shader.h | 5 ++++ .../drivers/radeonsi/si_state_draw.cpp | 6 +++++ 7 files changed, 37 insertions(+), 12 deletions(-) diff --git a/src/amd/common/ac_shader_util.c b/src/amd/common/ac_shader_util.c index 26a4d4c887b..173e9ed3672 100644 --- a/src/amd/common/ac_shader_util.c +++ b/src/amd/common/ac_shader_util.c @@ -1356,7 +1356,7 @@ ac_ngg_compute_subgroup_info(enum amd_gfx_level gfx_level, gl_shader_stage es_st enum mesa_prim input_prim, unsigned gs_vertices_out, unsigned gs_invocations, unsigned max_workgroup_size, unsigned wave_size, unsigned esgs_vertex_stride, unsigned ngg_lds_vertex_size, unsigned ngg_lds_scratch_size, bool tess_turns_off_ngg, - ac_ngg_subgroup_info *out) + unsigned max_esgs_lds_padding, ac_ngg_subgroup_info *out) { const unsigned gs_num_invocations = MAX2(gs_invocations, 1); const bool use_adjacency = mesa_prim_has_adjacency(input_prim); @@ -1365,7 +1365,7 @@ ac_ngg_compute_subgroup_info(enum amd_gfx_level gfx_level, gl_shader_stage es_st /* All these are in dwords. The maximum is 16K dwords (64KB) of LDS per workgroup. */ /* The LDS scratch is at the beginning of LDS space. */ - const unsigned max_lds_size = 16 * 1024 - ngg_lds_scratch_size / 4; + const unsigned max_lds_size = 16 * 1024 - ngg_lds_scratch_size / 4 - max_esgs_lds_padding / 4; const unsigned target_lds_size = max_lds_size; unsigned esvert_lds_size = 0; unsigned gsprim_lds_size = 0; diff --git a/src/amd/common/ac_shader_util.h b/src/amd/common/ac_shader_util.h index 184b26848d4..902fad5f5a2 100644 --- a/src/amd/common/ac_shader_util.h +++ b/src/amd/common/ac_shader_util.h @@ -344,7 +344,7 @@ ac_ngg_compute_subgroup_info(enum amd_gfx_level gfx_level, gl_shader_stage es_st enum mesa_prim input_prim, unsigned gs_vertices_out, unsigned gs_invocations, unsigned max_workgroup_size, unsigned wave_size, unsigned esgs_vertex_stride, unsigned ngg_lds_vertex_size, unsigned ngg_lds_scratch_size, bool tess_turns_off_ngg, - ac_ngg_subgroup_info *out); + unsigned max_esgs_lds_padding, ac_ngg_subgroup_info *out); #ifdef __cplusplus } diff --git a/src/amd/vulkan/radv_shader_info.c b/src/amd/vulkan/radv_shader_info.c index 90f2fb9190f..5dad01ef5f7 100644 --- a/src/amd/vulkan/radv_shader_info.c +++ b/src/amd/vulkan/radv_shader_info.c @@ -1339,7 +1339,7 @@ gfx10_get_ngg_info(const struct radv_device *device, struct radv_shader_info *es ac_ngg_compute_subgroup_info(gfx_level, es_info->stage, !!gs_info, input_prim, gs_vertices_out, gs_num_invocations, 128, stage_info->wave_size, es_info->esgs_itemsize, stage_info->ngg_lds_vertex_size, - stage_info->ngg_lds_scratch_size, false, &info); + stage_info->ngg_lds_scratch_size, false, 0, &info); out->hw_max_esverts = info.hw_max_esverts; out->max_gsprims = info.max_gsprims; diff --git a/src/gallium/drivers/radeonsi/si_nir_lower_abi.c b/src/gallium/drivers/radeonsi/si_nir_lower_abi.c index 8023c9c7a03..02549f8a0ab 100644 --- a/src/gallium/drivers/radeonsi/si_nir_lower_abi.c +++ b/src/gallium/drivers/radeonsi/si_nir_lower_abi.c @@ -606,6 +606,9 @@ static bool lower_intrinsic(nir_builder *b, nir_instr *instr, struct lower_abi_s case nir_intrinsic_load_polygon_stipple_buffer_amd: replacement = si_nir_load_internal_binding(b, args, SI_PS_CONST_POLY_STIPPLE, 4); break; + case nir_intrinsic_load_lds_ngg_gs_out_vertex_base_amd: + replacement = nir_imul_imm(b, GET_FIELD_NIR(GS_STATE_GS_OUT_LDS_OFFSET_256B), 256); + break; default: return false; } diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 9d71dd09f64..ffd8f1b24e9 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -196,15 +196,14 @@ static bool si_shader_binary_open(struct si_screen *screen, struct si_shader *sh (sel->stage <= MESA_SHADER_GEOMETRY && shader->key.ge.as_ngg))) { struct ac_rtld_symbol *sym = &lds_symbols[num_lds_symbols++]; sym->name = "esgs_ring"; - sym->size = (shader->key.ge.as_ngg ? shader->ngg.info.esgs_lds_size - : shader->gs_info.esgs_lds_size) * 4; + sym->size = 0; sym->align = 64 * 1024; } if (sel->stage == MESA_SHADER_GEOMETRY && shader->key.ge.as_ngg) { struct ac_rtld_symbol *sym = &lds_symbols[num_lds_symbols++]; sym->name = "ngg_emit"; - sym->size = shader->ngg.info.ngg_out_lds_size * 4; + sym->size = 0; sym->align = 4; } @@ -464,6 +463,10 @@ static void calculate_needed_lds_size(struct si_screen *sscreen, struct si_shade shader->config.lds_size = DIV_ROUND_UP(size_in_dw * 4, get_lds_granularity(sscreen, stage)); } + + /* Check that the LDS size is within hw limits. */ + assert(shader->config.lds_size * get_lds_granularity(sscreen, stage) <= + (sscreen->info.gfx_level == GFX6 ? 32 : 64) * 1024); } static int upload_binary_raw(struct si_screen *sscreen, struct si_shader *shader, @@ -518,8 +521,6 @@ static int upload_binary_raw(struct si_screen *sscreen, struct si_shader *shader post_upload_binary(sscreen, shader, rx_ptr, code_size, code_size, dma_upload, upload_ctx, staging, staging_offset); - - calculate_needed_lds_size(sscreen, shader); return code_size; } @@ -529,13 +530,17 @@ int si_shader_binary_upload_at(struct si_screen *sscreen, struct si_shader *shad bool dma_upload = !(sscreen->debug_flags & DBG(NO_DMA_SHADERS)) && sscreen->info.has_cp_dma && sscreen->info.has_dedicated_vram && !sscreen->info.all_vram_visible && bo_offset < 0; + int r; if (shader->binary.type == SI_SHADER_BINARY_ELF) { - return upload_binary_elf(sscreen, shader, scratch_va, dma_upload, bo_offset); + r = upload_binary_elf(sscreen, shader, scratch_va, dma_upload, bo_offset); } else { assert(shader->binary.type == SI_SHADER_BINARY_RAW); - return upload_binary_raw(sscreen, shader, scratch_va, dma_upload, bo_offset); + r = upload_binary_raw(sscreen, shader, scratch_va, dma_upload, bo_offset); } + + calculate_needed_lds_size(sscreen, shader); + return r; } int si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader, @@ -2583,10 +2588,16 @@ bool si_create_shader_variant(struct si_screen *sscreen, struct ac_llvm_compiler si_get_max_workgroup_size(shader), shader->wave_size, es_sel->info.esgs_vertex_stride, shader->info.ngg_lds_vertex_size, shader->info.ngg_lds_scratch_size, gs_sel->tess_turns_off_ngg, - &shader->ngg.info)) { + gs_sel->stage == MESA_SHADER_GEOMETRY ? 255 : 0, &shader->ngg.info)) { mesa_loge("Failed to compute subgroup info"); return false; } + + /* GS outputs in LDS must start at a multiple of 256B because GS_STATE_GS_OUT_LDS_OFFSET_256B + * doesn't store the low 8 bits. + */ + if (sel->stage == MESA_SHADER_GEOMETRY) + shader->ngg.info.esgs_lds_size = align(shader->ngg.info.esgs_lds_size, 64); /* align to 256B in dword units */ } else if (sscreen->info.gfx_level >= GFX9 && sel->stage == MESA_SHADER_GEOMETRY) { ac_legacy_gs_compute_subgroup_info(sel->info.base.gs.input_primitive, sel->info.base.gs.vertices_out, diff --git a/src/gallium/drivers/radeonsi/si_shader.h b/src/gallium/drivers/radeonsi/si_shader.h index a1a85c52d6a..4a2cb6900b4 100644 --- a/src/gallium/drivers/radeonsi/si_shader.h +++ b/src/gallium/drivers/radeonsi/si_shader.h @@ -238,6 +238,11 @@ enum * in the shader via vs_state_bits in legacy GS, the GS copy shader, and any NGG shader. */ /* bit gap */ +/* The LDS size of ES outputs in bytes for NGG GS, in multiples of 256 (bits [8:15]). + * This is used to determine the LDS address of GS outputs, which is after ES outputs. + */ +#define GS_STATE_GS_OUT_LDS_OFFSET_256B__SHIFT 6 +#define GS_STATE_GS_OUT_LDS_OFFSET_256B__MASK 0xff /* The number of ES outputs is derived from the last output index of SI_UNIQUE_SLOT_* + 1, which * can be 55 at most. The ESGS vertex stride in dwords is: NUM_ES_OUTPUTS * 4 + 1 * Only used by GFX9+ to compute LDS addresses of GS inputs. diff --git a/src/gallium/drivers/radeonsi/si_state_draw.cpp b/src/gallium/drivers/radeonsi/si_state_draw.cpp index 029a1ca601f..bd91e23b239 100644 --- a/src/gallium/drivers/radeonsi/si_state_draw.cpp +++ b/src/gallium/drivers/radeonsi/si_state_draw.cpp @@ -154,6 +154,12 @@ static bool si_update_shaders(struct si_context *sctx) } } + if (NGG && HAS_GS) { + assert((sctx->shader.gs.current->ngg.info.esgs_lds_size * 4) % 256 == 0); + SET_FIELD(sctx->current_gs_state, GS_STATE_GS_OUT_LDS_OFFSET_256B, + ((uint32_t)sctx->shader.gs.current->ngg.info.esgs_lds_size * 4) >> 8); + } + struct si_shader *api_vs = si_get_api_vs_inline(sctx, GFX_VERSION, HAS_TESS, HAS_GS); sctx->vs_uses_base_instance = api_vs->info.uses_base_instance; sctx->vs_uses_draw_id = api_vs->info.uses_draw_id;