comaniac commented on a change in pull request #7616:
URL: https://github.com/apache/tvm/pull/7616#discussion_r595362038
##########
File path: cmake/config.cmake
##########
@@ -99,6 +99,9 @@ set(USE_STACKVM_RUNTIME OFF)
# Whether enable tiny embedded graph runtime.
set(USE_GRAPH_RUNTIME ON)
+# Whether enable tiny graph runtime for cudaGraph Launch
Review comment:
```suggestion
# Whether enable tiny graph runtime with cudaGraph
```
##########
File path: cmake/modules/CUDA.cmake
##########
@@ -65,6 +65,17 @@ if(USE_CUDA)
list(APPEND RUNTIME_SRCS ${CONTRIB_THRUST_SRC})
endif(USE_THRUST)
+ if(USE_GRAPH_RUNTIME_CUGRAPH)
+ if(NOT USE_GRAPH_RUNTIME)
+ message(FATAL_ERROR "CUDA Graph is only supported by graph runtime,
should set USE_GRAPH_RUNTIME=ON")
Review comment:
```suggestion
message(FATAL_ERROR "CUDA Graph is only supported by graph runtime,
please set USE_GRAPH_RUNTIME=ON")
```
##########
File path: tests/python/unittest/test_runtime_graph_cugraph.py
##########
@@ -0,0 +1,92 @@
+# 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 os
+import re
+import sys
+import time
+
+import pytest
+
+import tvm
+import tvm.testing
+from tvm import te
+import numpy as np
+
+from tvm.contrib import utils, graph_runtime
+from tvm.contrib.cu_graph import cugraph_runtime
+
+
+bx = te.thread_axis("blockIdx.x")
+tx = te.thread_axis("threadIdx.x")
+
+
[email protected]_cuda
+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")
+ ctx = tvm.gpu(0)
+ try:
+ mod = cugraph_runtime.create(graph, mlib, ctx)
Review comment:
No worries. The added interface looks good.
##########
File path: tests/python/unittest/test_runtime_graph_cugraph.py
##########
@@ -0,0 +1,92 @@
+# 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 os
+import re
+import sys
+import time
+
+import pytest
+
+import tvm
+import tvm.testing
+from tvm import te
+import numpy as np
+
+from tvm.contrib import utils, graph_runtime
+from tvm.contrib.cu_graph import cugraph_runtime
+
+
+bx = te.thread_axis("blockIdx.x")
+tx = te.thread_axis("threadIdx.x")
+
+
[email protected]_cuda
+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")
+ ctx = tvm.gpu(0)
+ try:
+ mod = cugraph_runtime.create(graph, mlib, ctx)
+ except ValueError:
+ return
+ mod.capture_cuda_graph()
+ a = np.random.uniform(size=(n,)).astype(A.dtype)
+ mod.set_input(x=a)
+ mod.run_cuda_graph()
Review comment:
Make sense.
##########
File path: cmake/modules/CUDA.cmake
##########
@@ -65,6 +65,17 @@ if(USE_CUDA)
list(APPEND RUNTIME_SRCS ${CONTRIB_THRUST_SRC})
endif(USE_THRUST)
+ if(USE_GRAPH_RUNTIME_CUGRAPH)
+ if(NOT USE_GRAPH_RUNTIME)
+ message(FATAL_ERROR "CUDA Graph is only supported by graph runtime,
should set USE_GRAPH_RUNTIME=ON")
+ endif()
+ if(CUDAToolkit_VERSION_MAJOR LESS "10")
+ message(FATAL_ERROR "CUDA Graph requires at least CUDA 10, got="
${CUDAToolkit_VERSION})
+ endif()
+ message(STATUS "Build with Graph runtime cuGraph support...")
Review comment:
It would be better to have one terminology in this PR. Either `cuGraph`
or `CUDA graph`.
##########
File path: python/tvm/testing.py
##########
@@ -514,6 +514,23 @@ def requires_cuda(*args):
return _compose(args, _requires_cuda)
+def requires_cudagraph(*args):
+ """Mark a test as requiring the CUDA Graph Feature
+
+ This also marks the test as requiring cuda
+
+ Parameters
+ ----------
+ f : function
+ Function to mark
+ """
+ _requires_cudagraph = [
+ pytest.mark.skipif(not nvcc.have_cudagraph(), reason="CUDA Graph not
support"),
Review comment:
```suggestion
pytest.mark.skipif(not nvcc.have_cudagraph(), reason="CUDA Graph is
not supported in this environment"),
```
##########
File path: cmake/modules/CUDA.cmake
##########
@@ -65,6 +65,17 @@ if(USE_CUDA)
list(APPEND RUNTIME_SRCS ${CONTRIB_THRUST_SRC})
endif(USE_THRUST)
+ if(USE_GRAPH_RUNTIME_CUGRAPH)
+ if(NOT USE_GRAPH_RUNTIME)
+ message(FATAL_ERROR "CUDA Graph is only supported by graph runtime,
should set USE_GRAPH_RUNTIME=ON")
+ endif()
+ if(CUDAToolkit_VERSION_MAJOR LESS "10")
+ message(FATAL_ERROR "CUDA Graph requires at least CUDA 10, got="
${CUDAToolkit_VERSION})
Review comment:
```suggestion
message(FATAL_ERROR "CUDA Graph requires CUDA 10 or above, got="
${CUDAToolkit_VERSION})
```
##########
File path: src/runtime/graph/graph_runtime_factory.cc
##########
@@ -130,6 +138,31 @@ Module GraphRuntimeFactory::DebugRuntimeCreate(const
std::vector<TVMContext>& ct
return mod;
}
+Module GraphRuntimeFactory::CudaGraphRuntimeCreate(const
std::vector<TVMContext>& ctxs) {
+ const PackedFunc* pf =
tvm::runtime::Registry::Get("tvm.graph_runtime_cugraph.create");
+ ICHECK(pf != nullptr) << "Cannot find function
tvm.graph_runtime_cugraph.create in registry. "
+ "Do you enable cuda graph runtime build?";
Review comment:
```suggestion
"Did you set(USE_GRAPH_RUNTIME_CUGRAPH=ON)?";
```
##########
File path: python/tvm/contrib/nvcc.py
##########
@@ -349,6 +349,19 @@ def have_tensorcore(compute_version=None, target=None):
return False
+def have_cudagraph():
+ """Either CUDA Graph support is provided"""
+ try:
+ cuda_path = find_cuda_path()
+ cuda_ver = get_cuda_version(cuda_path)
+ if cuda_ver < 10.0:
+ return False
+ return True
+ except RuntimeError:
+ warnings.warn("Cannot find cuda path")
Review comment:
This warning has no information and can consider to remove.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
[email protected]
