You're on the right track. Generally, the way to run something on Hexagon is
to run an app on the CPU, and have it offload code to Hexagon via the FastRPC
mechanism. If your Hexagon code has function `foo`, and you want to call it
from the CPU, you create the IDL description of `foo`s interface, and the
generated stub/skel libraries are the CPU/Hexagon sides respectively of the
remote call to `foo`.
The direct communication via C or C++ program running on CPU is certainly
possible, but it's not necessarily convenient for use with TVM: you would use
TVM to build a shared library for Hexagon, then you would need to write that
C/C++ code that would load it and execute it.
TVM does have a framework to make the whole thing easier. There is a tool
called "RPC tracker", which serves as a hub to which actual RPC servers
connect. An app wanting to offload something would connect to the tracker, and
from the tracker it would get the connection with the corresponding RPC server.
The app could then upload the code for remote execution and run it.
There is already an app that performs the role of the RPC server: it's the
cpp_rpc program, so most of the pieces of the puzzle are present.
Now, the workflow is basically as follows:
1. Start the RPC tracker on a PC. It will typically listen on port 9190. Once
it starts, leave it running.
2. Set up port forwarding on the device (this is so that the TCP/IP connection
between the tracker and the app will work):
```
adb forward tcp:5001 tcp:5001
adb reverse tcp:9190 tcp:9190
```
3. Copy the tvm_rpc app to some directory on the device. I usually copy the
stub libraries there as well (but it may work with the libraries in /vendor/lib
as well). Then start the app (the key is just some identifier that the RPC
server will register itself under with the tracker):
```
cd /the/path/on/device
export LD_LIBRARY_PATH=.:${LD_LIBRARY_PATH} # this is if stub libs are in
this dir
export ADSP_LIBRARY_PATH=/vendor/lib/rfsa/adsp:${ADSP_LIBRARY_PATH}
./tvm_rpc server --host=127.0.0.1 --port=5001 --tracker=127.0.0.1:9190
--key=<some-word>
```
4. At this point you can run a TVM python code on your PC and have it
communicate with the device.
>From the point of view of the python code itself, you need to indicate in your
>schedule which part of your kernel needs to be outlined. If `s` is the
>schedule, and `Op` is the operation which you want to offload, mark it as
>`pipeline`:
```
px, x = s[Op].split(s[Op].op.axis[0], nparts=1)
s[Op].bind(px, tvm.te.thread_axis("pipeline"))
```
Then do `tvm.build` with
```
target = tvm.target.hexagon()
target_host = llvm -mtriple=aarch64-unknown-linux-android26
```
(you can use a different version of Android as the OS here).
The output of `tvm.build`, let's call it `m`, will be a module with code built
for AArch64. It will also contain another module in it with the code for
Hexagon. Now you should be able to save both of these modules individually:
```
m.save('cpu.so')
m.imported_modules[0].save('hexagon.so') # this will also create a .json file
```
At this point you can establish a connection:
```
tracker = tvm.rpc.connect_tracker(tracker_host, tracker_port)
remote = tracker.request(key, priority=0, session_timeout=60)
```
The `tracket_host`/`tracker_port` are the hostname and port number of the
tracker. You can use `127.0.0.1` and `9190` here. The `key` parameter is the
identifier you used with `tvm_rpc` app.
Finally, upload all the files to the device, and them load them on the device
(you may need to add full path names):
```
remote.upload('cpu.so')
remote.upload('hexagon.so')
remote.upload('hexagon.tvm_meta.json')
remote_cpu = remote.load_module('cpu.so')
remote_hexagon = remote.load_module('hexagon.so')
remote_cpu.import_module(remote_hexagon)
```
Here, `remote_cpu` is the callable remote module which you can invoke:
```
remote_cpu(a)
```
Remember that `a` has to have the remote Hexagon context, i.e. `ctx =
remote.hexagon(0)`.
Feel free to ask more questions if something isn't clear.
---
[Visit
Topic](https://discuss.tvm.apache.org/t/implementation-of-hexagon-runtime-for-target/9284/2)
to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click
here](https://discuss.tvm.apache.org/email/unsubscribe/cc798e30eda12727d6e54ff728897fac7fe2fd1cf11309296307013e803e3c7e).
