mesa/src/broadcom
Jose Maria Casanova Crespo d95076e581 v3dv: lower oversized compute workgroups to 256 invocations
V3D advertises maxComputeWorkGroupInvocations = 256 but ggml-vulkan
in many cases ignores this limit an creates compute pipelines with
over this limit. Although this is a bug in the application we can
take advantage of nir_lower_workgroup_size and make the application
work.

This issue was causing an assertion failure at nir_to_vir.c:

  assert(c->local_invocation_index_bits <= 8);

The solution is lowering the oversized workgroups to a 256-invocation
workgroup loop, like radv and radeonsi are doing on GFX7, by running
nir_lower_workgroup_size(256) for this scenario.

Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41257>
2026-04-30 13:59:19 +00:00
..
ci ci: Delete references to various broken traces. 2026-04-22 17:39:31 +00:00
cle broadcom/cle: parse once the XML spec 2026-04-06 09:17:15 +00:00
clif broadcom: use Mesa logging functions 2026-04-06 07:40:55 +00:00
common v3d/v3dv: Use new V3D_MAX_CSD_WG_SIZE = 256 2026-04-30 13:59:18 +00:00
compiler broadcom/compiler: move nir_lower_undef_to_zero out of optimization loop 2026-04-30 12:30:34 +02:00
drm-shim broadcom: use Mesa logging functions 2026-04-06 07:40:55 +00:00
ds mesa: replace most occurrences of getenv() with os_get_option() 2025-11-06 04:36:13 +00:00
perfcntrs broadcom: use Mesa logging functions 2026-04-06 07:40:55 +00:00
qpu broadcom/compiler: Add V3D 7.1 v8dot dot product QPU instructions 2026-04-29 13:21:07 +00:00
simulator broadcom: use Mesa logging functions 2026-04-06 07:40:55 +00:00
vulkan v3dv: lower oversized compute workgroups to 256 invocations 2026-04-30 13:59:19 +00:00
.editorconfig
meson.build broadcom: Add perfetto data source 2024-12-13 12:29:11 +00:00