Reviewed-by: Bas Nieuwenhuizen <b...@basnieuwenhuizen.nl>
On Fri, Feb 1, 2019 at 12:07 PM Timothy Arceri <tarc...@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@lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
