fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-06-03 15:28:15 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 4

kk: Fix pre-compiled compute grid size

Browse source 
Vulkan dispatch commands should multiply group count by local size,
but pre-compiled dispatches should not. For example, the predicate
indirect shader has a local size of 32 and a grid size equal to the
max draw count, which was resulting in a dispatch total grid size of
(max_draw_count * 32), overrunning the buffer.

Reviewed-by: Aitor Camacho <aitor@lunarg.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41740>
This commit is contained in:
squidbus 2026-05-21 18:25:30 -07:00 • committed by Marge Bot
parent b8f0fe6bdc
commit 22f61a4eb5
2 changed files with 7 additions and 9 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

7
src/kosmickrisp/bridge/mtl_encoder.m
View file

@ -208,11 +208,8 @@ mtl_dispatch_threads(mtl_compute_encoder *encoder,
{
@autoreleasepool {
id<MTLComputeCommandEncoder> enc = (id<MTLComputeCommandEncoder>)encoder;
MTLSize thread_count = MTLSizeMake(grid_size.x * local_size.x,
grid_size.y * local_size.y,
grid_size.z * local_size.z);
MTLSize threads_per_threadgroup = MTLSizeMake(local_size.x,
local_size.y,
MTLSize thread_count = MTLSizeMake(grid_size.x, grid_size.y, grid_size.z);
MTLSize threads_per_threadgroup = MTLSizeMake(local_size.x, local_size.y,
local_size.z);
// TODO_KOSMICKRISP can we rely on nonuniform threadgroup size support?

9
src/kosmickrisp/vulkan/kk_cmd_dispatch.c
View file

@ -70,13 +70,14 @@ kk_CmdDispatchBase(VkCommandBuffer commandBuffer, uint32_t baseGroupX,
kk_flush_compute_state(cmd);
struct kk_shader *cs = cmd->state.shaders[MESA_SHADER_COMPUTE];
struct mtl_size local_size = cs->info.cs.local_size;
struct mtl_size grid_size = {
.x = groupCountX,
.y = groupCountY,
.z = groupCountZ,
.x = groupCountX * local_size.x,
.y = groupCountY * local_size.y,
.z = groupCountZ * local_size.z,
};
mtl_compute_encoder *enc = kk_compute_encoder(cmd);
mtl_dispatch_threads(enc, grid_size, cs->info.cs.local_size);
mtl_dispatch_threads(enc, grid_size, local_size);
}
VKAPI_ATTR void VKAPI_CALL