radeonsi: stop using LLVM LDS linking logic for the GS out LDS offset

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 <pierre-eric.pelloux-prayer@amd.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35529>

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;

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
 }

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;

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;
    }

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,

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.

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;