apache / apache/tvm

[Bug] model cross compiled with cutlass can't run

Open
#19,422 0 comments 0 reactions 0 assignees View on GitHub
needs-triage type: bug
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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.