ITADN

Metal: set_vector_bytes passes nullptr to setBytes for empty vectors

#4061OpenJamesPriceZV 创建于 11 天前
buglow priority
J
JamesPriceZVcommented
### 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 条评论