[Mesa-dev] [PATCH] radv: take LDS into account for compute shader occupancy stats

[Bas Nieuwenhuizen]

Reviewed-by: Bas Nieuwenhuizen <bas at basnieuwenhuizen.nl>

On Fri, Feb 1, 2019 at 12:07 PM Timothy Arceri <tarceri at itsqueeze.com> wrote:
>
> Ported from d205faeb6c96.
> ---
>  src/amd/vulkan/radv_nir_to_llvm.c |  6 +++---
>  src/amd/vulkan/radv_private.h     |  3 +++
>  src/amd/vulkan/radv_shader.c      | 10 ++++++++--
>  3 files changed, 14 insertions(+), 5 deletions(-)
>
> diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c
> index e80938527e5..d90a4c0de1e 100644
> --- a/src/amd/vulkan/radv_nir_to_llvm.c
> +++ b/src/amd/vulkan/radv_nir_to_llvm.c
> @@ -3372,9 +3372,9 @@ ac_setup_rings(struct radv_shader_context *ctx)
>         }
>  }
>
> -static unsigned
> -ac_nir_get_max_workgroup_size(enum chip_class chip_class,
> -                             const struct nir_shader *nir)
> +unsigned
> +radv_nir_get_max_workgroup_size(enum chip_class chip_class,
> +                               const struct nir_shader *nir)
>  {
>         switch (nir->info.stage) {
>         case MESA_SHADER_TESS_CTRL:
> diff --git a/src/amd/vulkan/radv_private.h b/src/amd/vulkan/radv_private.h
> index 85c18906f84..e5b8286ea62 100644
> --- a/src/amd/vulkan/radv_private.h
> +++ b/src/amd/vulkan/radv_private.h
> @@ -1934,6 +1934,9 @@ void radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm,
>                              int nir_count,
>                              const struct radv_nir_compiler_options *options);
>
> +unsigned radv_nir_get_max_workgroup_size(enum chip_class chip_class,
> +                                        const struct nir_shader *nir);
> +
>  /* radv_shader_info.h */
>  struct radv_shader_info;
>
> diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c
> index 07450ff236b..a7fce02ee83 100644
> --- a/src/amd/vulkan/radv_shader.c
> +++ b/src/amd/vulkan/radv_shader.c
> @@ -744,7 +744,8 @@ generate_shader_stats(struct radv_device *device,
>                       gl_shader_stage stage,
>                       struct _mesa_string_buffer *buf)
>  {
> -       unsigned lds_increment = device->physical_device->rad_info.chip_class >= CIK ? 512 : 256;
> +       enum chip_class chip_class = device->physical_device->rad_info.chip_class;
> +       unsigned lds_increment = chip_class >= CIK ? 512 : 256;
>         struct ac_shader_config *conf;
>         unsigned max_simd_waves;
>         unsigned lds_per_wave = 0;
> @@ -757,12 +758,17 @@ generate_shader_stats(struct radv_device *device,
>                 lds_per_wave = conf->lds_size * lds_increment +
>                                align(variant->info.fs.num_interp * 48,
>                                      lds_increment);
> +       } else if (stage == MESA_SHADER_COMPUTE) {
> +               unsigned max_workgroup_size =
> +                               ac_nir_get_max_workgroup_size(chip_class, variant->nir);
> +               lds_per_wave = (conf->lds_size * lds_increment) /
> +                              DIV_ROUND_UP(max_workgroup_size, 64);
>         }
>
>         if (conf->num_sgprs)
>                 max_simd_waves =
>                         MIN2(max_simd_waves,
> -                            ac_get_num_physical_sgprs(device->physical_device->rad_info.chip_class) / conf->num_sgprs);
> +                            ac_get_num_physical_sgprs(chip_class) / conf->num_sgprs);
>
>         if (conf->num_vgprs)
>                 max_simd_waves =
> --
> 2.20.1
>
> _______________________________________________
> mesa-dev mailing list
> mesa-dev at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev