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;