tree 3ba41a8b34ca76dda2a4d3d9ff4874e3345a7d31
parent 954f90d2f53b22ff8b347ae215d58b6cb1ca1cb3
author Alex Vasile <48962821+Alex-Vasile@users.noreply.github.com> 1782727222 -0400
committer GitHub <noreply@github.com> 1782727222 +0200
gpgsig -----BEGIN PGP SIGNATURE-----
 
 wsFcBAABCAAQBQJqQkI2CRC1aQ7uu5UhlAAAkzMQABd1CkuLopJzCJ6FeNH05zIN
 lUPXla9ZLSwSfLVPLIHvGW6U7X0tFsi6QuK35NZpf38G+DvkP7IGrwM0+uTECx46
 w14YhDXeH4l/9ne7TNBRb8f1lSifBcrP1qQDuPNQX6olD+oOtAdRYz32Qh9O46Oj
 DqSe0S0bKeKQfugAw5v5vOasTWYDFWUIa3YSph3Hb7bpN6oV2LTpPvrkrZaYFXU6
 zREw5fvXyjBCqrOFnACDQpDVjkSVjzqxR2YzBVlMLsoXRGKV/NOokKM8tftf5bQ1
 IaL7RCNwV5Xsu+OWZU5WoU4ZN2TTmQ0RpxjXM5pGOmqbzlAw6782odVBJNFSo7Ac
 gIxYEXHZDOSLSuG9DDd6xrqEiQmdV8+fKY1Zl7qOl9pd6y2lWhmsuL2jqxDH2NMn
 UGj/OK9yzfaSRWJZ20dZjAgj5Tv5WpoUCl3dKvMLX/377YsIVDSBfUMk8VlfvQeE
 ZTghGYfrJve8lte7SiHNieTI4aMatyjL/x9Vlmgbz3JE9pLQQIth80nTqSmtnbZm
 6s9LIrB8nfyOz6XR9efN2FIIsnKMgnr2+mA3XF8NALd9275r0jRN98Zrjb5B+s89
 JGG9K4Pq3liVmABQ/7m4ESwY0/MneAmBZFSlLh7RrV7UOiUA44NXTE++1I82UU5Y
 oNm87u6caN8LfvQMfk75
 =fNBU
 -----END PGP SIGNATURE-----
 

[Metal] Fix copy_buffer_1byte dispatch grid (#24640)

copy_buffer_1byte silently dropped the trailing ~3/4 of every copy that
fell back to the compute-kernel path. The kernel copies one byte per
thread, but its dispatch grid was sized as ceil(length / (workgroup_size
* 4)) -- a quarter of the required ceil(length / workgroup_size)
workgroups -- so only the first min(ceil(length/128)*32, length) bytes
were written and the tail was left untouched. The bounds guard `if (id
>= spec.length) return` prevents over-run but cannot launch the missing
threads.

On macOS the compute path is taken whenever the source offset, target
offset, or length is not a multiple of 4 (the Metal blit fast-path
requires all three 4-byte aligned). It is also the path behind the
staging->device upload
(iree_hal_metal_command_buffer_prepare_update_buffer), so non-4-aligned
source uploads were truncated too -- a device copy then read a zeroed
source tail.

The *4 multiplier is correct for the fill kernels (fill_buffer_16byte /
fill_buffer_4byte / fill_buffer_1byte process 16/4/4 bytes per thread
respectively), which is where it was copied from; copy_buffer_1byte is a
flat 1-byte-per-thread loop with no such expansion.

Fix (dispatch only -- the copy_buffer_1byte.metal kernel is unchanged):
size the grid as ceil(length / workgroup_size) so the launched thread
count (ceil(length/32) * 32) covers every byte. The existing `if (id >=
spec.length) return` guard masks the surplus threads in the final
workgroup.

Verified:
CTS/CommandBufferCopyBufferTest.CopySizeAndAlignmentClasses/metal now
passes all alignment/size sub-cases that previously truncated
(offset/length not a multiple of 4; lengths 31/32/33/.../65536 across
all alignment classes). The aligned16_mib sub-case still fails,
independently, from a staging-buffer capacity issue (128 KiB staging
buffer vs. a 1 MiB single-command-buffer upload), not from this dispatch
sizing. Signed-off-by: Alex Vasile
<48962821+Alex-Vasile@users.noreply.github.com>

Signed-off-by: Alex Vasile <48962821+Alex-Vasile@users.noreply.github.com>