diff options
| author | Timothy Arceri <[email protected]> | 2019-02-01 22:04:39 +1100 |
|---|---|---|
| committer | Timothy Arceri <[email protected]> | 2019-02-01 22:25:30 +1100 |
| commit | 9b9ccee4d64b5e64f6638bca7a87b3f159e3be9c (patch) | |
| tree | f3f336bd3556aa435a0abd77f07313ee8891bc08 /src/amd | |
| parent | a53d68d3185f1eee4eccc5ddc8d8519c62df3686 (diff) | |
radv: take LDS into account for compute shader occupancy stats
Ported from d205faeb6c96.
Reviewed-by: Bas Nieuwenhuizen <[email protected]>
Diffstat (limited to 'src/amd')
| -rw-r--r-- | src/amd/vulkan/radv_nir_to_llvm.c | 6 | ||||
| -rw-r--r-- | src/amd/vulkan/radv_private.h | 3 | ||||
| -rw-r--r-- | 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 2fd287fcd17..3e5edcc54a7 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 = |