radv: move more CS info to gather_shader_info_cs()
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com> Reviewed-by: Timur Kristóf <timur.kristof@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18210>
This commit is contained in:

committed by
Marge Bot

parent
ac47d149c6
commit
cae4eb2904
@@ -3420,34 +3420,6 @@ radv_fill_shader_info(struct radv_pipeline *pipeline,
|
|||||||
|
|
||||||
radv_nir_shader_info_link(device, pipeline_key, stages, last_vgt_api_stage);
|
radv_nir_shader_info_link(device, pipeline_key, stages, last_vgt_api_stage);
|
||||||
|
|
||||||
if (stages[MESA_SHADER_COMPUTE].nir) {
|
|
||||||
unsigned subgroup_size = pipeline_key->cs.compute_subgroup_size;
|
|
||||||
unsigned req_subgroup_size = subgroup_size;
|
|
||||||
bool require_full_subgroups = pipeline_key->cs.require_full_subgroups;
|
|
||||||
|
|
||||||
if (!subgroup_size)
|
|
||||||
subgroup_size = device->physical_device->cs_wave_size;
|
|
||||||
|
|
||||||
unsigned local_size = stages[MESA_SHADER_COMPUTE].nir->info.workgroup_size[0] *
|
|
||||||
stages[MESA_SHADER_COMPUTE].nir->info.workgroup_size[1] *
|
|
||||||
stages[MESA_SHADER_COMPUTE].nir->info.workgroup_size[2];
|
|
||||||
|
|
||||||
/* Games don't always request full subgroups when they should,
|
|
||||||
* which can cause bugs if cswave32 is enabled.
|
|
||||||
*/
|
|
||||||
if (device->physical_device->cs_wave_size == 32 &&
|
|
||||||
stages[MESA_SHADER_COMPUTE].nir->info.cs.uses_wide_subgroup_intrinsics && !req_subgroup_size &&
|
|
||||||
local_size % RADV_SUBGROUP_SIZE == 0)
|
|
||||||
require_full_subgroups = true;
|
|
||||||
|
|
||||||
if (require_full_subgroups && !req_subgroup_size) {
|
|
||||||
/* don't use wave32 pretending to be wave64 */
|
|
||||||
subgroup_size = RADV_SUBGROUP_SIZE;
|
|
||||||
}
|
|
||||||
|
|
||||||
stages[MESA_SHADER_COMPUTE].info.cs.subgroup_size = subgroup_size;
|
|
||||||
}
|
|
||||||
|
|
||||||
for (int i = 0; i < MESA_VULKAN_SHADER_STAGES; i++) {
|
for (int i = 0; i < MESA_VULKAN_SHADER_STAGES; i++) {
|
||||||
if (stages[i].nir) {
|
if (stages[i].nir) {
|
||||||
stages[i].info.wave_size = radv_get_wave_size(device, i, &stages[i].info);
|
stages[i].info.wave_size = radv_get_wave_size(device, i, &stages[i].info);
|
||||||
|
@@ -506,9 +506,34 @@ gather_shader_info_fs(const nir_shader *nir, const struct radv_pipeline_key *pip
|
|||||||
}
|
}
|
||||||
|
|
||||||
static void
|
static void
|
||||||
gather_shader_info_cs(const nir_shader *nir, struct radv_shader_info *info)
|
gather_shader_info_cs(struct radv_device *device, const nir_shader *nir,
|
||||||
|
const struct radv_pipeline_key *pipeline_key, struct radv_shader_info *info)
|
||||||
{
|
{
|
||||||
info->cs.uses_ray_launch_size = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_RAY_LAUNCH_SIZE_ADDR_AMD);
|
info->cs.uses_ray_launch_size = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_RAY_LAUNCH_SIZE_ADDR_AMD);
|
||||||
|
|
||||||
|
unsigned subgroup_size = pipeline_key->cs.compute_subgroup_size;
|
||||||
|
unsigned req_subgroup_size = subgroup_size;
|
||||||
|
bool require_full_subgroups = pipeline_key->cs.require_full_subgroups;
|
||||||
|
|
||||||
|
if (!subgroup_size)
|
||||||
|
subgroup_size = device->physical_device->cs_wave_size;
|
||||||
|
|
||||||
|
unsigned local_size =
|
||||||
|
nir->info.workgroup_size[0] * nir->info.workgroup_size[1] * nir->info.workgroup_size[2];
|
||||||
|
|
||||||
|
/* Games don't always request full subgroups when they should, which can cause bugs if cswave32
|
||||||
|
* is enabled.
|
||||||
|
*/
|
||||||
|
if (device->physical_device->cs_wave_size == 32 && nir->info.cs.uses_wide_subgroup_intrinsics &&
|
||||||
|
!req_subgroup_size && local_size % RADV_SUBGROUP_SIZE == 0)
|
||||||
|
require_full_subgroups = true;
|
||||||
|
|
||||||
|
if (require_full_subgroups && !req_subgroup_size) {
|
||||||
|
/* don't use wave32 pretending to be wave64 */
|
||||||
|
subgroup_size = RADV_SUBGROUP_SIZE;
|
||||||
|
}
|
||||||
|
|
||||||
|
info->cs.subgroup_size = subgroup_size;
|
||||||
}
|
}
|
||||||
|
|
||||||
static void
|
static void
|
||||||
@@ -648,7 +673,7 @@ radv_nir_shader_info_pass(struct radv_device *device, const struct nir_shader *n
|
|||||||
|
|
||||||
switch (nir->info.stage) {
|
switch (nir->info.stage) {
|
||||||
case MESA_SHADER_COMPUTE:
|
case MESA_SHADER_COMPUTE:
|
||||||
gather_shader_info_cs(nir, info);
|
gather_shader_info_cs(device, nir, pipeline_key, info);
|
||||||
break;
|
break;
|
||||||
case MESA_SHADER_TASK:
|
case MESA_SHADER_TASK:
|
||||||
gather_shader_info_task(nir, info);
|
gather_shader_info_task(nir, info);
|
||||||
|
Reference in New Issue
Block a user