[Bug] model cross compiled with cutlass can't run
- Dominant language
- Python
- Stars
- 13.7k
- Forks
- 4k
- Avg merge
- 2d 1h
- Merged PRs (30d)
- 112
Description
Hi everyone,
Currently I’m using tvm 0.23.0 ver with cutlass trying to deploy resnet model on a nvidia thor device. But I get an error.
### Expected behavior
model running with cutlass backend
### Actual behavior
The log error shows here
```
RPC succ
start upload
lib upload ok
start load_module
remote load_module ok
Traceback (most recent call last):
File "/data/opensource/tvm/tvm_tutorial/test03_e2e/tuning_onnx_01/deploy_thorU_rev11_question.py", line 136, in
run()
File "/data/opensource/tvm/tvm_tutorial/test03_e2e/tuning_onnx_01/deploy_thorU_rev11_question.py", line 118, in run
vm = relax.VirtualMachine(lib, [dev])
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/opensource/tvm/tvm/python/tvm/runtime/vm.py", line 91, in __init__
self._setup_device(device, memory_cfg)
File "/data/opensource/tvm/tvm/python/tvm/runtime/vm.py", line 124, in _setup_device
self.module["vm_initialization"](*init_args)
File "python/tvm_ffi/cython/function.pxi", line 923, in tvm_ffi.core.Function.__call__
File "/data/opensource/tvm/tvm/src/runtime/rpc/rpc_module.cc", line 145, in tvm::runtime::RPCWrappedFunc::operator()(tvm::ffi::PackedArgs, tvm::ffi::Any*) const
sess_->CallFunc(handle_, ffi::PackedArgs(packed_args.data(), packed_args.size()), set_return);
File "/data/opensource/tvm/tvm/src/runtime/rpc/rpc_endpoint.cc", line 1116, in tvm::runtime::RPCClientSession::CallFunc(void*, tvm::ffi::PackedArgs, std::function const&)
endpoint_->CallFunc(func, args, fencode_return);
File "/data/opensource/tvm/tvm/src/runtime/rpc/rpc_endpoint.cc", line 906, in tvm::runtime::RPCEndpoint::CallFunc(void*, tvm::ffi::PackedArgs, std::function)
code = HandleUntilReturnEvent(true, encode_return);
File "/data/opensource/tvm/tvm/src/runtime/rpc/rpc_endpoint.cc", line 746, in tvm::runtime::RPCEndpoint::HandleUntilReturnEvent(bool, std::function)
code = handler_->HandleNextEvent(client_mode, false, setreturn);
File "/data/opensource/tvm/tvm/src/runtime/rpc/rpc_endpoint.cc", line 134, in tvm::runtime::RPCEndpoint::EventHandler::HandleNextEvent(bool, bool, std::function)
this->HandleProcessPacket(setreturn);
File "/data/opensource/tvm/tvm/src/runtime/rpc/rpc_endpoint.cc", line 409, in tvm::runtime::RPCEndpoint::EventHandler::HandleProcessPacket(std::function)
this->HandleReturn(code, setreturn);
File "/data/opensource/tvm/tvm/src/runtime/rpc/rpc_endpoint.cc", line 473, in tvm::runtime::RPCEndpoint::EventHandler::HandleReturn(tvm::runtime::RPCCode, std::function)
LOG(FATAL) << msg;
File "/data/opensource/tvm/tvm/include/tvm/runtime/logging.h", line 321, in tvm::runtime::detail::LogFatal::~LogFatal()
GetEntry().Finalize();
File "/data/opensource/tvm/tvm/include/tvm/runtime/logging.h", line 337, in tvm::runtime::detail::LogFatal::Entry::Finalize()
InternalError error(file_, lineno_, stream_.str());
tvm.error.RPCError: Error caught from RPC call:
Check failed: (func.has_value()) is false: Error: Cannot find ffi::Function fused_relax_nn_conv2d_cutlass in either Relax VM kernel library, or in TVM runtime ffi::Function registry, or in global Relax functions of the VM executable
```
Any environment details, such as: Operating System, TVM version, etc
tvm 0.23.0
### Steps to reproduce
Preferably a minimal script to cause the issue to occur.
```
import os
import tvm
from tvm import relax
from tvm.relax.frontend.onnx import from_onnx
from tvm.relax.transform import ConvertLayout
import tvm.relax.backend.cuda.cutlass as relax_cutlass
from tvm.contrib import cutlass as cutlass_utils
import onnx
import numpy as np
ONNX_PATH = "resnet18.onnx"
LIB_PATH = "/tmp/model/model_deployed.so"
LIB = "model_deployed.so"
TMP_DIR = "/tmp/model/build"
TRACKER_HOST = "127.0.0.1"
TRACKER_PORT = 9090
SM = 101
TARGET_HOST = tvm.target.Target("llvm -mtriple=aarch64-linux-gnu")
TARGET = tvm.target.Target("cuda -arch=sm_101", host=TARGET_HOST)
def _patch_cutlass_cross_compile(sm: int):
import tvm.contrib.cutlass.build as _cb
_orig = _cb._get_cutlass_compile_options
def _patched(sm, threads, use_fast_math=False):
kwargs = _orig(sm, threads, use_fast_math)
kwargs["options"] += [
"--compiler-bindir", "aarch64-linux-gnu-g++",
]
print(f"[patch] nvcc options: {kwargs['options']}")
return kwargs
_cb._get_cutlass_compile_options = _patched
print(f"[patch] _get_cutlass_compile_options replaced")
def build_with_cutlass(mod):
_patch_cutlass_cross_compile(SM)
with TARGET:
mod = relax.transform.FoldConstant()(mod)
mod = relax.transform.LegalizeOps(skip_ops=["relax.nn.conv2d"])(mod)
mod = ConvertLayout({"relax.nn.conv2d": ["NHWC", "OHWI"]})(mod)
mod = relax_cutlass.partition_for_cutlass(mod)
print(f"[CUTLASS] part num: {cutlass_utils.num_cutlass_partitions(mod)}")
mod = relax.transform.RunCodegen(
{"cutlass": {"sm": SM, "find_first_valid": False}}
)(mod)
mod = relax.transform.LegalizeOps()(mod)
mod = relax.transform.LambdaLift()(mod)
from tvm import dlight as dl
mod = dl.ApplyDefaultSchedule(
dl.gpu.Matmul(), dl.gpu.GEMV(),
dl.gpu.Reduction(), dl.gpu.Fallback(),
)(mod)
executable = relax.build(mod, target=TARGET)
return executable
def export_library_cross_so(executable, output_so, tmp_dir):
import os
os.makedirs(tmp_dir, exist_ok=True)
os.makedirs(os.path.dirname(output_so), exist_ok=True)
executable.export_library(
output_so,
cc="aarch64-linux-gnu-g++",
workspace_dir=tmp_dir,
)
print(f"[export] export done: {output_so}")
return output_so
def run():
os.makedirs(TMP_DIR, exist_ok=True)
onnx_model = onnx.load(ONNX_PATH)
mod = from_onnx(onnx_model, keep_params_in_input=False)
executable = build_with_cutlass(mod)
export_library_cross_so(executable, LIB_PATH, TMP_DIR)
tracker = tvm.rpc.connect_tracker(url=TRACKER_HOST, port=TRACKER_PORT)
remote = tracker.request(key="thor", session_timeout=600)
print("RPC succ")
print("start upload")
remote.upload(LIB_PATH)
print("lib upload ok")
print("start load_module")
lib = remote.load_module(LIB)
print("remote load_module ok")
dev = remote.cuda(0)
vm = relax.VirtualMachine(lib, [dev])
print("create vm ok")
input_np = np.random.rand(1, 3, 224, 224).astype("float32")
remote_input = tvm.runtime.tensor(input_np, device=dev)
vm.set_input("main", remote_input)
print("set input ok")
vm.invoke_stateful("main")
print("call invoke stateful ok")
time_f = vm.time_evaluator("invoke_stateful", dev, number=10, repeat=3)
prof = time_f("main")
print(f"[perf] {prof.mean * 1000:.2f} ± {prof.std * 1000:.2f} ms")
if __name__ == "__main__":
try:
run()
except Exception:
import traceback
traceback.print_exc()
```
so My Question is:
1. Is my script is the correct way using cutlass optimizing a model? Please recommend me one if I am wrong.
2. why that error happens? please give me some hints to solve it, thank you.
Contributor guide
No contributing guide indexed for this repository
Research direction
Start by running the provided minimal script and compare its failure at relax.VirtualMachine(lib, [dev]) with the missing fused_relax_nn_conv2d_cutlass function. Read tvm/runtime/vm.py and the referenced RPC runtime files, then trace the Cutlass partition and RunCodegen entry points in the script. Done means identifying why the generated Cutlass function is unavailable on the remote VM and confirming a working cross-compiled model run.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python
- Domain
- compilers, machine-learning
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Needs clarification
- Newbie friendliness
- 35/100