From: Timothy Arceri <tarceri@itsqueeze.com>
Date: Fri, 1 Feb 2019 11:04:39 +0000 (+1100)
Subject: radv: take LDS into account for compute shader occupancy stats
X-Git-Tag: upstream/19.3.0~10095
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=9b9ccee4d64b5e64f6638bca7a87b3f159e3be9c;p=platform%2Fupstream%2Fmesa.git

radv: take LDS into account for compute shader occupancy stats

Ported from d205faeb6c96.

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

diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c
index e809385..d90a4c0 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 85c1890..e5b8286 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 2fd287f..3e5edcc 100644
--- a/src/amd/vulkan/radv_shader.c
+++ b/src/amd/vulkan/radv_shader.c
@@ -733,7 +733,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;
@@ -746,12 +747,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 =