tqchen commented on issue #578:
URL: https://github.com/apache/tvm-ffi/issues/578#issuecomment-4361102776
Thanks for the note. This actually needs a bit more thoughts. (2) is acually
not that complicated, since `DLTensor.data` should actualy be `id<MTLBuffer>`
(i believe it already is since for opaque cases like metal data should be the
buffer object)
The main issue is (1), because ideally we want some reasonably ABI agnostic
way of passing the current command queue for MPS, and ideally we also want it
to be reasonably generic (e.g. ideally if things can work for mlx if possible).
Unlike cuda, metal do not have the concept of default device and command queue,
so commandqueue handling is somewhat done per framework. There is also not a
common interface that we can rely on on job submission. At the momentum, there
are different approaches in exposing things (e.g. does things eager vs lazy,
how much command encoder reuse across kernels). One thing that could work is
some form of generic interface and have an addon `tvm-ffi-metal` that allows
registration of such interface, so torch variant can be passed and have the
kernels depend on the addon, and when the respective torch extension get
imported the default interface set to use torch's variant(while allowing other
variants e.g. if it is desirable to lazily reuse command encoder w
hich i believe is what mlx is doing for host perf reasons).
```c++
class TVMFFIMetalStreamInterface {
public:
virtual id<MTLComputeCommandEncoder> AcquireEncoder() = 0;
virtual void ReleaseEncoder() = 0;
virtual id<MTLDevice> GetDevice() = 0;
virtual dispatch_queue_t GetDispatchQueue() = 0;
virtual ~TVMFFIMetalStream() = default;
};
```
The above code can be a sketch of how this can looks like (but likely need
more thoughts to really be confident that it is the stable low-level interface
that can work). Taking a step back, one can of course directly call
`at::mps::getCurrentMPSStream()` in tvm ffi defined interface, but of course
that brings a direct dependency on libtorch.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
