Skip to content

[BugFix][Metal] Reuse borrowed stream wrappers and size device slots correctly - #76

Open
anerli wants to merge 1 commit into
tile-ai:tilelang_mainfrom
magnitudedev:pr/metal-borrowed-streams
Open

anerli wants to merge 1 commit into
tile-ai:tilelang_mainfrom
magnitudedev:pr/metal-borrowed-streams

Conversation

@anerli

@anerli anerli commented Sep 15, 2026

Copy link
Copy Markdown

Problem and root cause

metal.SetStream allocates a new MetalRawStream wrapper on every call and overwrites the previous pointer without freeing the wrapper. Its growth path also resizes the stream vector to device_id before writing element device_id, leaving that element outside the vector's size.

Change

  • src/runtime/metal/metal_common.h: allow an existing borrowed-stream wrapper to update its command-buffer pointer.
  • src/runtime/metal/metal_module.mm: own one wrapper per thread/device with a thread-local vector of unique_ptrs, reuse it on subsequent calls, and grow device slots to device_id + 1.

The wrapper borrows the command buffer; the change owns the wrapper rather than taking ownership of the external command buffer.

Tests and validation

A local native Objective-C++ regression harness linked against this branch's rebuilt TVM objects passes checks for:

  • Growth from an empty stream vector to device slot 0.
  • Stable wrapper identity across 10,000 calls, with replacement command-buffer pointers.
  • Growth to slot 3 while preserving slot 0 and keeping distinct wrappers per device.
  • Distinct wrappers in separate threads.

On unmodified upstream, the harness fails the first slot-size assertion. The harness is local validation, not an automated regression test included in this branch. It uses sentinel command-buffer pointers without dispatching GPU work.

pre-commit run --files src/runtime/metal/metal_common.h \
  src/runtime/metal/metal_module.mm

The existing Metal execution suite also has a baseline limitation: test_metal_inf_nan segfaults on unmodified upstream. The native checks above validate wrapper bookkeeping, not end-to-end borrowed-stream GPU execution.

Validation environment: Apple M4 Max, macOS/Darwin 24.5.0 arm64, Python 3.12.11. Each branch was rebuilt separately from upstream 907a88c87 with:

cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release \
  -DUSE_METAL=ON -DUSE_LLVM=OFF -DUSE_CUDA=OFF -DUSE_Z3=OFF
cmake --build build --target tvm tvm_compiler tvm_runtime -j 8

git diff --check passes. The applicable formatting and other pre-commit hooks pass. The repository-wide ASF-header hook reports 11 existing violations, reproduced on unmodified upstream; none is in the changed files.

Related

No duplicate open PR was found. This runtime fix is independent of Metal target-property discovery in #71.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant