)]}'
{
  "commit": "51bd22fe7d6e3cf72692d367b84d7c3b4eb8f35b",
  "tree": "3ba41a8b34ca76dda2a4d3d9ff4874e3345a7d31",
  "parents": [
    "954f90d2f53b22ff8b347ae215d58b6cb1ca1cb3"
  ],
  "author": {
    "name": "Alex Vasile",
    "email": "48962821+Alex-Vasile@users.noreply.github.com",
    "time": "Mon Jun 29 06:00:22 2026 -0400"
  },
  "committer": {
    "name": "GitHub",
    "email": "noreply@github.com",
    "time": "Mon Jun 29 12:00:22 2026 +0200"
  },
  "message": "[Metal] Fix copy_buffer_1byte dispatch grid (#24640)\n\ncopy_buffer_1byte silently dropped the trailing ~3/4 of every copy that\nfell back to the compute-kernel path. The kernel copies one byte per\nthread, but its dispatch grid was sized as ceil(length / (workgroup_size\n* 4)) -- a quarter of the required ceil(length / workgroup_size)\nworkgroups -- so only the first min(ceil(length/128)*32, length) bytes\nwere written and the tail was left untouched. The bounds guard `if (id\n\u003e\u003d spec.length) return` prevents over-run but cannot launch the missing\nthreads.\n\nOn macOS the compute path is taken whenever the source offset, target\noffset, or length is not a multiple of 4 (the Metal blit fast-path\nrequires all three 4-byte aligned). It is also the path behind the\nstaging-\u003edevice upload\n(iree_hal_metal_command_buffer_prepare_update_buffer), so non-4-aligned\nsource uploads were truncated too -- a device copy then read a zeroed\nsource tail.\n\nThe *4 multiplier is correct for the fill kernels (fill_buffer_16byte /\nfill_buffer_4byte / fill_buffer_1byte process 16/4/4 bytes per thread\nrespectively), which is where it was copied from; copy_buffer_1byte is a\nflat 1-byte-per-thread loop with no such expansion.\n\nFix (dispatch only -- the copy_buffer_1byte.metal kernel is unchanged):\nsize the grid as ceil(length / workgroup_size) so the launched thread\ncount (ceil(length/32) * 32) covers every byte. The existing `if (id \u003e\u003d\nspec.length) return` guard masks the surplus threads in the final\nworkgroup.\n\nVerified:\nCTS/CommandBufferCopyBufferTest.CopySizeAndAlignmentClasses/metal now\npasses all alignment/size sub-cases that previously truncated\n(offset/length not a multiple of 4; lengths 31/32/33/.../65536 across\nall alignment classes). The aligned16_mib sub-case still fails,\nindependently, from a staging-buffer capacity issue (128 KiB staging\nbuffer vs. a 1 MiB single-command-buffer upload), not from this dispatch\nsizing. Signed-off-by: Alex Vasile\n\u003c48962821+Alex-Vasile@users.noreply.github.com\u003e\n\nSigned-off-by: Alex Vasile \u003c48962821+Alex-Vasile@users.noreply.github.com\u003e",
  "tree_diff": [
    {
      "type": "modify",
      "old_id": "a14ece35c79163f55867e4c4424130c109de4cec",
      "old_mode": 33188,
      "old_path": "runtime/src/iree/hal/drivers/metal/builtin_executables.m",
      "new_id": "8ed1cc4601e2641dc2e801417038dff3e4b741cb",
      "new_mode": 33188,
      "new_path": "runtime/src/iree/hal/drivers/metal/builtin_executables.m"
    }
  ]
}