Metal: set_vector_bytes passes nullptr to setBytes for empty vectors
buglow priority
### Describe the bug
`CommandEncoder::set_vector_bytes` forwards `vec.data()` straight to `MTLComputeCommandEncoder::setBytes`. For an empty vector `data()` is permitted to return `nullptr`, and `setBytes(nullptr, 0, idx)` is rejected by Metal. With Metal API Validation enabled the process aborts:
```
-[MTLDebugComputeCommandEncoder setBytes:length:atIndex:]:
failed assertion 'bytes argument cannot be nil'
```
Current code in `mlx/backend/metal/device.h`:
```cpp
template <typename Vec, typename = std::enable_if_t<is_vector_v<Vec>>>
void set_vector_bytes(const Vec& vec, size_t nelems, int idx) {
get_command_encoder()->setBytes(
vec.data(), nelems * sizeof(typename Vec::value_type), idx);
}
```
Several indexing specializations legitimately pass an empty metadata vector whose binding the kernel never reads, so reaching this path is normal rather than a caller error.
### To Reproduce
Run any indexing/gather path that binds empty shape or stride metadata, with Metal API Validation enabled (Xcode scheme diagnostics, or `MTL_DEBUG_LAYER=1`). We hit it during `TokenIterator.prepare` on a 4-bit Gemma model through mlx-swift.
Observed stack:
```
mlx::core::Gather::eval_gpu
-> CommandEncoder::set_vector_bytes
-> -[MTLDebugComputeCommandEncoder setBytes:length:atIndex:]
-> __assert_rtn / abort
```
Note this is undefined behaviour regardless of validation; the debug layer only makes it visible.
### Expected behavior
Binding an empty, unread vector should not abort. Binding a single zeroed element preserves the kernel contract and satisfies Metal:
```cpp
template <typename Vec, typename = std::enable_if_t<is_vector_v<Vec>>>
void set_vector_bytes(const Vec& vec, size_t nelems, int idx) {
using Value = typename Vec::value_type;
if (nelems == 0) {
static const Value empty_value{};
get_command_encoder()->setBytes(&empty_value, sizeof(Value), idx);
return;
}
get_command_encoder()->setBytes(vec.data(), nelems * sizeof(Value), idx);
}
```
Guarding at the call sites instead would also work, but there are several and the invariant is easy to reintroduce.
### Desktop (please complete the following information)
- OS Version: macOS 27.0
- Version: reproduced against `mlx-swift` 0.31.6 (MLX core `ce45c525`); the code path is unchanged on current `main` (`47bbfe8`)
2 条评论