From 54045780aaa1f2b1815c0e518415c2f6d31586c8 Mon Sep 17 00:00:00 2001 From: Sviatoslav Peleshko Date: Thu, 31 Mar 2022 16:44:15 +0300 Subject: anv: workaround apps that assume full subgroups without specifying it MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Without this we might choose 8 or 16 width, while the app assumes 32. With subgroup operations it may cause wrong calculations and thus bugs. Examples of such games are Aperture Desk Job and DOOM Eternal. v2: Make it a driconf option instead of applying unconditionally, move from brw_required_dispatch_width to brw_compile_cs v3: Rename allow_assuming_full_subgroups -> assume_full_subgroups. Include assume_full_subgroups value in anv_pipeline_hash_compute(). v4: Move actual workaround code from brw_fs.c -> anv_pipeline.c. Cc: mesa-stable Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/6171 Signed-off-by: Sviatoslav Peleshko Reviewed-by: Lionel Landwerlin Reviewed-by: Marcin Ślusarz Part-of: (cherry picked from commit 28ca5636f6519f70cede02742f5ba0e00e6afcd3) --- .pick_status.json | 2 +- src/intel/vulkan/anv_device.c | 4 ++++ src/intel/vulkan/anv_pipeline.c | 35 ++++++++++++++++++++++++++++------- src/intel/vulkan/anv_private.h | 5 +++++ src/util/00-mesa-defaults.conf | 8 ++++++++ src/util/driconf.h | 8 ++++++++ 6 files changed, 54 insertions(+), 8 deletions(-) diff --git a/.pick_status.json b/.pick_status.json index c06d1df55a0..52ab9072931 100644 --- a/.pick_status.json +++ b/.pick_status.json @@ -236,7 +236,7 @@ "description": "anv: workaround apps that assume full subgroups without specifying it", "nominated": true, "nomination_type": 0, - "resolution": 0, + "resolution": 1, "because_sha": null }, { diff --git a/src/intel/vulkan/anv_device.c b/src/intel/vulkan/anv_device.c index 25db66fb607..b39e9b2f4bb 100644 --- a/src/intel/vulkan/anv_device.c +++ b/src/intel/vulkan/anv_device.c @@ -68,6 +68,7 @@ static const driOptionDescription anv_dri_options[] = { DRI_CONF_VK_X11_OVERRIDE_MIN_IMAGE_COUNT(0) DRI_CONF_VK_X11_STRICT_IMAGE_COUNT(false) DRI_CONF_VK_XWAYLAND_WAIT_READY(true) + DRI_CONF_ANV_ASSUME_FULL_SUBGROUPS(false) DRI_CONF_SECTION_END DRI_CONF_SECTION_DEBUG @@ -1093,6 +1094,9 @@ anv_init_dri_options(struct anv_instance *instance) instance->vk.app_info.app_version, instance->vk.app_info.engine_name, instance->vk.app_info.engine_version); + + instance->assume_full_subgroups = + driQueryOptionb(&instance->dri_options, "anv_assume_full_subgroups"); } VkResult anv_CreateInstance( diff --git a/src/intel/vulkan/anv_pipeline.c b/src/intel/vulkan/anv_pipeline.c index 94774e82cdd..2ffb4fc32a9 100644 --- a/src/intel/vulkan/anv_pipeline.c +++ b/src/intel/vulkan/anv_pipeline.c @@ -738,9 +738,14 @@ anv_pipeline_hash_compute(struct anv_compute_pipeline *pipeline, if (layout) _mesa_sha1_update(&ctx, layout->sha1, sizeof(layout->sha1)); - const bool rba = pipeline->base.device->robust_buffer_access; + const struct anv_device *device = pipeline->base.device; + + const bool rba = device->robust_buffer_access; _mesa_sha1_update(&ctx, &rba, sizeof(rba)); + const bool afs = device->physical->instance->assume_full_subgroups; + _mesa_sha1_update(&ctx, &afs, sizeof(afs)); + _mesa_sha1_update(&ctx, stage->shader_sha1, sizeof(stage->shader_sha1)); _mesa_sha1_update(&ctx, &stage->key.cs, sizeof(stage->key.cs)); @@ -2061,7 +2066,8 @@ anv_pipeline_compile_cs(struct anv_compute_pipeline *pipeline, }; int64_t pipeline_start = os_time_get_nano(); - const struct brw_compiler *compiler = pipeline->base.device->physical->compiler; + struct anv_device *device = pipeline->base.device; + const struct brw_compiler *compiler = device->physical->compiler; struct anv_pipeline_stage stage = { .stage = MESA_SHADER_COMPUTE, @@ -2090,8 +2096,8 @@ anv_pipeline_compile_cs(struct anv_compute_pipeline *pipeline, const enum brw_subgroup_size_type subgroup_size_type = anv_subgroup_size_type(MESA_SHADER_COMPUTE, stage.module, info->stage.flags, rss_info); - populate_cs_prog_key(&pipeline->base.device->info, subgroup_size_type, - pipeline->base.device->robust_buffer_access, + populate_cs_prog_key(&device->info, subgroup_size_type, + device->robust_buffer_access, &stage.key.cs); ANV_FROM_HANDLE(anv_pipeline_layout, layout, info->layout); @@ -2103,7 +2109,7 @@ anv_pipeline_compile_cs(struct anv_compute_pipeline *pipeline, bool cache_hit = false; if (!skip_cache_lookup) { - bin = anv_device_search_for_kernel(pipeline->base.device, cache, + bin = anv_device_search_for_kernel(device, cache, &stage.cache_key, sizeof(stage.cache_key), &cache_hit); @@ -2138,6 +2144,21 @@ anv_pipeline_compile_cs(struct anv_compute_pipeline *pipeline, anv_pipeline_lower_nir(&pipeline->base, mem_ctx, &stage, layout); + unsigned local_size = stage.nir->info.workgroup_size[0] * + stage.nir->info.workgroup_size[1] * + stage.nir->info.workgroup_size[2]; + + /* Games don't always request full subgroups when they should, + * which can cause bugs, as they may expect bigger size of the + * subgroup than we choose for the execution. + */ + if (device->physical->instance->assume_full_subgroups && + stage.nir->info.cs.uses_wide_subgroup_intrinsics && + subgroup_size_type == BRW_SUBGROUP_SIZE_API_CONSTANT && + local_size && + local_size % BRW_SUBGROUP_SIZE == 0) + stage.key.base.subgroup_size_type = BRW_SUBGROUP_SIZE_REQUIRE_32; + stage.num_stats = 1; struct brw_compile_cs_params params = { @@ -2145,7 +2166,7 @@ anv_pipeline_compile_cs(struct anv_compute_pipeline *pipeline, .key = &stage.key.cs, .prog_data = &stage.prog_data.cs, .stats = stage.stats, - .log_data = pipeline->base.device, + .log_data = device, }; stage.code = brw_compile_cs(compiler, mem_ctx, ¶ms); @@ -2163,7 +2184,7 @@ anv_pipeline_compile_cs(struct anv_compute_pipeline *pipeline, } const unsigned code_size = stage.prog_data.base.program_size; - bin = anv_device_upload_kernel(pipeline->base.device, cache, + bin = anv_device_upload_kernel(device, cache, MESA_SHADER_COMPUTE, &stage.cache_key, sizeof(stage.cache_key), stage.code, code_size, diff --git a/src/intel/vulkan/anv_private.h b/src/intel/vulkan/anv_private.h index b8e16d17199..c1644f3e7a1 100644 --- a/src/intel/vulkan/anv_private.h +++ b/src/intel/vulkan/anv_private.h @@ -1051,6 +1051,11 @@ struct anv_instance { struct driOptionCache dri_options; struct driOptionCache available_dri_options; + + /** + * Workarounds for game bugs. + */ + bool assume_full_subgroups; }; VkResult anv_init_wsi(struct anv_physical_device *physical_device); diff --git a/src/util/00-mesa-defaults.conf b/src/util/00-mesa-defaults.conf index f110836961a..4e4576af8f6 100644 --- a/src/util/00-mesa-defaults.conf +++ b/src/util/00-mesa-defaults.conf @@ -885,6 +885,14 @@ TODO: document the other workarounds.