diff --git a/src/intel/vulkan/anv_device.c b/src/intel/vulkan/anv_device.c
index 5fc4d4567d0..c67c57bec16 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
@@ -1100,6 +1101,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 2004246453d..1c37a572a44 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));
@@ -1915,7 +1920,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,
@@ -1944,8 +1950,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);
@@ -1957,7 +1963,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);
@@ -1992,6 +1998,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 = {
@@ -1999,7 +2020,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);
@@ -2017,7 +2038,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 8c3a34f2788..fa4c8662a5c 100644
--- a/src/intel/vulkan/anv_private.h
+++ b/src/intel/vulkan/anv_private.h
@@ -1079,6 +1079,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 03cbdc5af67..40d53ca040c 100644
--- a/src/util/00-mesa-defaults.conf
+++ b/src/util/00-mesa-defaults.conf
@@ -914,6 +914,14 @@ TODO: document the other workarounds.
+
+
+
+
+
+
+
+
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