Mesa (staging/22.1): anv: workaround apps that assume full subgroups without specifying it
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Tue Apr 26 16:22:28 UTC 2022
Module: Mesa
Branch: staging/22.1
Commit: fe83982fa62a7fa6a19a5f8728cdcf9732cfd5f9
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=fe83982fa62a7fa6a19a5f8728cdcf9732cfd5f9
Author: Sviatoslav Peleshko <sviatoslav.peleshko at globallogic.com>
Date: Thu Mar 31 16:44:15 2022 +0300
anv: workaround apps that assume full subgroups without specifying it
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 <sviatoslav.peleshko at globallogic.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin at intel.com>
Reviewed-by: Marcin Ślusarz <marcin.slusarz at intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15708>
(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 2d4039c58bb..7c68883a4e6 100644
--- a/.pick_status.json
+++ b/.pick_status.json
@@ -256,7 +256,7 @@
"description": "anv: workaround apps that assume full subgroups without specifying it",
"nominated": true,
"nomination_type": 0,
- "resolution": 0,
+ "resolution": 1,
"main_sha": null,
"because_sha": null
},
diff --git a/src/intel/vulkan/anv_device.c b/src/intel/vulkan/anv_device.c
index 1438a5dcf3f..b52858e1057 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
@@ -1097,6 +1098,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 eb5da15cfff..7a365b0f155 100644
--- a/src/intel/vulkan/anv_pipeline.c
+++ b/src/intel/vulkan/anv_pipeline.c
@@ -640,9 +640,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));
@@ -1914,7 +1919,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,
@@ -1943,8 +1949,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);
@@ -1956,7 +1962,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);
@@ -1991,6 +1997,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 = {
@@ -1998,7 +2019,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);
@@ -2016,7 +2037,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 caa463f6958..4df086fa650 100644
--- a/src/intel/vulkan/anv_private.h
+++ b/src/intel/vulkan/anv_private.h
@@ -1082,6 +1082,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 59de17e2110..917904a09a9 100644
--- a/src/util/00-mesa-defaults.conf
+++ b/src/util/00-mesa-defaults.conf
@@ -910,6 +910,14 @@ TODO: document the other workarounds.
<option name="vs_position_always_invariant" value="true" />
</application>
</device>
+ <device driver="anv">
+ <application name="Aperture Desk Job" executable="deskjob">
+ <option name="anv_assume_full_subgroups" value="true" />
+ </application>
+ <application name="DOOMEternal" executable="DOOMEternalx64vk.exe">
+ <option name="anv_assume_full_subgroups" value="true" />
+ </application>
+ </device>
<device driver="virtio_gpu">
<application name="Counter-Strike Global Offensive" executable="csgo_linux64">
diff --git a/src/util/driconf.h b/src/util/driconf.h
index fc720374554..e580f6fe2d8 100644
--- a/src/util/driconf.h
+++ b/src/util/driconf.h
@@ -576,4 +576,12 @@
DRI_CONF_OPT_B(radv_disable_aniso_single_level, def, \
"Disable anisotropic filtering for single level images")
+/**
+ * \brief ANV specific configuration options
+ */
+
+#define DRI_CONF_ANV_ASSUME_FULL_SUBGROUPS(def) \
+ DRI_CONF_OPT_B(anv_assume_full_subgroups, def, \
+ "Allow assuming full subgroups requirement even when it's not specified explicitly")
+
#endif
More information about the mesa-commit
mailing list