This is an automated email from the ASF dual-hosted git repository.
syfeng pushed a commit to branch refactor
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/refactor by this push:
new 1936a83f0e fix
1936a83f0e is described below
commit 1936a83f0e886876129cc30714104b924d6fa161
Author: Siyuan Feng <[email protected]>
AuthorDate: Sat Feb 15 17:37:50 2025 +0800
fix
---
python/tvm/contrib/cuda_graph/__init__.py | 16 ---
.../tvm/contrib/cuda_graph/cuda_graph_executor.py | 134 ---------------------
rust/tvm-rt/Cargo.toml | 1 -
tests/python/relax/test_dataflow_pattern.py | 8 +-
.../runtime/test_runtime_graph_cuda_graph.py | 93 --------------
.../test_tir_transform_fp8_legalize.py | 4 -
6 files changed, 4 insertions(+), 252 deletions(-)
diff --git a/python/tvm/contrib/cuda_graph/__init__.py
b/python/tvm/contrib/cuda_graph/__init__.py
deleted file mode 100644
index 13a83393a9..0000000000
--- a/python/tvm/contrib/cuda_graph/__init__.py
+++ /dev/null
@@ -1,16 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
diff --git a/python/tvm/contrib/cuda_graph/cuda_graph_executor.py
b/python/tvm/contrib/cuda_graph/cuda_graph_executor.py
deleted file mode 100644
index d047316eb5..0000000000
--- a/python/tvm/contrib/cuda_graph/cuda_graph_executor.py
+++ /dev/null
@@ -1,134 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Graph executor with CUDA Graph"""
-import tvm._ffi
-
-from tvm._ffi.base import string_types
-from tvm.contrib import graph_executor
-
-
-def create(graph_json_str, libmod, device):
- """Create a runtime executor module given a graph and module.
-
- Parameters
- ----------
- graph_json_str : str
- The graph to be deployed in json format output by json graph.
- The graph can contain operator(tvm_op) that points to the name
- of PackedFunc in the libmod.
-
- libmod : tvm.runtime.Module
- The module of the corresponding function
-
- device : Device
- The device to deploy the module, only supports CUDA GPU
-
- Returns
- -------
- graph_module : GraphModuleCudaGraph
- CUDA graph executor module that can be used to execute the graph.
-
- Note
- ----
- See also
:py:class:`tvm.contrib.cuda_graph.cuda_graph_executor.GraphModuleCudaGraph`
- for examples to directly construct a GraphModuleCudaGraph from an exported
- relay compiled library.
- """
- assert isinstance(graph_json_str, string_types)
- try:
- dev, num_rpc_dev, device_type_id = graph_executor.get_device(libmod,
device)
- if num_rpc_dev == len(dev):
- fcreate =
dev[0]._rpc_sess.get_function("tvm.graph_executor_cuda_graph.create")
- else:
- fcreate =
tvm._ffi.get_global_func("tvm.graph_executor_cuda_graph.create")
- except ValueError:
- raise ValueError(
- "To enable CUDA graph support (experimental), please set "
- "'(USE_GRAPH_EXECUTOR_CUGRAPH ON)' in config.cmake and rebuild TVM"
- )
-
- return GraphModuleCudaGraph(fcreate(graph_json_str, libmod,
*device_type_id))
-
-
-class GraphModuleCudaGraph(graph_executor.GraphModule):
- """CUDA graph executor module.
-
- This is a CUDA graph executor wrapper over the TVM runtime.
- Runtime interfaces are wrapped with CUDA graph functionalities.
-
- Parameters
- ----------
- module : Module
- The internal tvm module that holds the actual graph functions.
- """
-
- def __init__(self, module):
- self._start_capture = module["start_capture"]
- self._end_capture = module["end_capture"]
- self._run_cuda_graph = module["run_cuda_graph"]
- self._cuda_graph_captured = False
- graph_executor.GraphModule.__init__(self, module)
-
- def capture_cuda_graph(self):
- """Capture a CUDA graph for tvm_op graph
-
- This should be called before run_cuda_graph() to capture and
- instantiate a CUDA graph instance.
- """
- self._run() # call cuModuleLoadData before cudaStream API
- self._start_capture()
- self._run()
- self._end_capture()
- self._cuda_graph_captured = True
-
- def run_cuda_graph(self):
- """Run the CUDA graph for tvm_op graph
-
- Run the captured CUDA graph instance instead of the
- for-loop kernel launch of default graph executor
- """
- self._run_cuda_graph()
-
- def run(self, **input_dict):
- """A run wrapper for graph capture / launch, user can just
- change default graph executor to cuda graph executor, and
- the first call will capture a cuda graph for future launch
-
- Parameters
- ----------
- input_dict: dict of str to NDArray
- List of input values to be feed to
- """
- if input_dict:
- self.set_input(**input_dict)
- if not self._cuda_graph_captured:
- self.capture_cuda_graph()
- else:
- self._run_cuda_graph()
-
- def debug_get_output(self, node, out):
- """Run graph up to node and get the output to out
-
- Parameters
- ----------
- node : int / str
- The node index or name
-
- out : NDArray
- The output array container
- """
- raise NotImplementedError("Please use debugger.debug_executor as
graph_executor instead.")
diff --git a/rust/tvm-rt/Cargo.toml b/rust/tvm-rt/Cargo.toml
index 24d9061a21..8f61b76c58 100644
--- a/rust/tvm-rt/Cargo.toml
+++ b/rust/tvm-rt/Cargo.toml
@@ -56,7 +56,6 @@ use-micro = ["tvm-sys/use-micro"]
use-install-dev = ["tvm-sys/use-install-dev"]
hide-private-symbols = ["tvm-sys/hide-private-symbols"]
use-fallback-stl-map = ["tvm-sys/use-fallback-stl-map"]
-use-ethosn = ["tvm-sys/use-ethosn"]
use-index-default-i64 = ["tvm-sys/use-index-default-i64"]
use-tf-tvmdsoop = ["tvm-sys/use-tf-tvmdsoop"]
use-byodt-posit = ["tvm-sys/use-byodt-posit"]
diff --git a/tests/python/relax/test_dataflow_pattern.py
b/tests/python/relax/test_dataflow_pattern.py
index 3aa316cff7..4b5da0d9e6 100644
--- a/tests/python/relax/test_dataflow_pattern.py
+++ b/tests/python/relax/test_dataflow_pattern.py
@@ -278,13 +278,13 @@ def test_extern_fn_pattern():
def test_op_attr():
x = rx.Var("x", R.Tensor("float32"))
y = rx.Var("y", R.Tensor("float32"))
- conv2d = rx.nn.conv2d(x, y, kernel_size=(3, 3))
+ conv2d = rx.op.nn.conv2d(x, y, strides=(3, 3))
xp = is_var("x")
yp = is_var("y")
# TODO(@yuchen): reenable the assert after figuring out why it fails
- # assert is_op("nn.conv2d")(xp, yp).has_attr({"kernel_size": [3,
3]}).match(conv2d)
- assert not is_op("nn.conv2d")(xp, yp).has_attr({"kernel_size": [4,
3]}).match(conv2d)
- assert not is_op("nn.conv2d")(xp, yp).has_attr({"kernel_size_": [3,
3]}).match(conv2d)
+ # assert is_op("nn.conv2d")(xp, yp).has_attr({"strides": [3,
3]}).match(conv2d)
+ assert not is_op("nn.conv2d")(xp, yp).has_attr({"strides": [4,
3]}).match(conv2d)
+ assert not is_op("nn.conv2d")(xp, yp).has_attr({"strides": [3,
3]}).match(conv2d)
def test_match_call_attr():
diff --git a/tests/python/runtime/test_runtime_graph_cuda_graph.py
b/tests/python/runtime/test_runtime_graph_cuda_graph.py
deleted file mode 100644
index 3013b346b5..0000000000
--- a/tests/python/runtime/test_runtime_graph_cuda_graph.py
+++ /dev/null
@@ -1,93 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-import json
-
-import tvm
-import tvm.testing
-from tvm import te
-import numpy as np
-
-from tvm.contrib.cuda_graph import cuda_graph_executor
-
-
-bx = te.thread_axis("blockIdx.x")
-tx = te.thread_axis("threadIdx.x")
-
-
[email protected]_cudagraph
-def test_graph_simple():
- n = 32
- A = te.placeholder((n,), name="A")
- B = te.compute(A.shape, lambda *i: A(*i) + 1.0, name="B")
- s = te.create_schedule(B.op)
- xo, xi = s[B].split(B.op.axis[0], factor=8)
- s[B].bind(xo, bx)
- s[B].bind(xi, tx)
-
- node0 = {"op": "null", "name": "x", "inputs": []}
- node1 = {
- "op": "tvm_op",
- "name": "add",
- "inputs": [[0, 0, 0]],
- "attrs": {"func_name": "myadd", "flatten_data": "1", "num_inputs":
"1", "num_outputs": "1"},
- }
- nodes = [node0, node1]
- arg_nodes = [0]
- node_row_ptr = [0, 1, 2]
- outputs = [[1, 0, 0]]
- shape = (n,)
- attrs = {
- "shape": ["list_shape", [shape, shape]],
- "dltype": ["list_str", ["float32", "float32"]],
- "storage_id": ["list_int", [0, 1]],
- }
- graph = {
- "nodes": nodes,
- "arg_nodes": arg_nodes,
- "node_row_ptr": node_row_ptr,
- "heads": outputs,
- "attrs": attrs,
- }
- graph = json.dumps(graph)
-
- def check_verify():
- mlib = tvm.build(s, [A, B], "cuda", name="myadd")
- dev = tvm.cuda(0)
- try:
- mod = cuda_graph_executor.create(graph, mlib, dev)
- except ValueError:
- return
-
- for i in range(3):
- a = np.random.uniform(size=(n,)).astype(A.dtype)
- mod.run(x=a) # The first run captured a CUDA graph
- out = mod.get_output(0, tvm.nd.empty((n,)))
- np.testing.assert_equal(out.numpy(), a + 1)
-
- # capture / run CUDA graph manually
- mod.capture_cuda_graph()
- a = np.random.uniform(size=(n,)).astype(A.dtype)
- mod.set_input(x=a)
- mod.run_cuda_graph()
- out = mod.get_output(0, tvm.nd.empty((n,)))
- np.testing.assert_equal(out.numpy(), a + 1)
-
- check_verify()
-
-
-if __name__ == "__main__":
- test_graph_simple()
diff --git a/tests/python/tir-transform/test_tir_transform_fp8_legalize.py
b/tests/python/tir-transform/test_tir_transform_fp8_legalize.py
index 62e7072479..e1f487c572 100644
--- a/tests/python/tir-transform/test_tir_transform_fp8_legalize.py
+++ b/tests/python/tir-transform/test_tir_transform_fp8_legalize.py
@@ -206,8 +206,6 @@ promote_dtype = tvm.testing.parameter("float16", "float32")
def test_fp8_compute_legalize(dtype, promote_dtype):
- if
tvm.contrib.nvcc.have_fp8(tvm.contrib.nvcc.get_target_compute_version()):
- return
target = Target("cuda")
before = BindTarget(target)(get_before(dtype))
expected = BindTarget(target)(get_after_compute_legalize(dtype,
promote_dtype))
@@ -219,8 +217,6 @@ def test_fp8_compute_legalize(dtype, promote_dtype):
def test_fp8_storage_legalize(dtype, promote_dtype):
- if
tvm.contrib.nvcc.have_fp8(tvm.contrib.nvcc.get_target_compute_version()):
- return
target = Target("cuda")
before = BindTarget(target)(get_after_compute_legalize(dtype,
promote_dtype))
after = tvm.tir.transform.FP8StorageLegalize()(before)
