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

[Timothy Arceri]

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