yuchaoli opened a new pull request #8065: URL: https://github.com/apache/tvm/pull/8065 This pull request adds sparse conv2d(1*1) support for auto_scheduler. I have implemented sparse conv2d with kernel_size=1, using the BSR sparse data format. I evaluate the performance of sparse conv2d by mobilenet_v2 in pixel2. All of 1*1 conv2d in the network are transformed into sparse conv2d with block_size=(4, 1). | Sparse Rate | 0 %| 10 %| 20 %| 30 %| 40 %| 50 %| 60 %| 70 %| 80 %| 90 %| |---------|-------|-------|-------|-------|-------|-------|-------|-------|-------|-------| | Latency | 17.37ms| 21.85ms | 21.05ms | 18.18ms | 17.16ms | 15.03ms | 14.06ms | 11.59ms | 10.12ms | 8.67ms | @jcf94 -- 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: users@infra.apache.org
yuchaoli commented on a change in pull request #8065:
URL: https://github.com/apache/tvm/pull/8065#discussion_r634060483
##########
File path: tests/python/relay/test_sparse_conv2d_convert.py
##########
@@ -0,0 +1,110 @@
+# 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 itertools
+
+import numpy as np
+import scipy.sparse as sp
+
+
+import tvm
+from tvm.ir import IRModule
+from tvm import relay
+
+
+def random_bsr_matrix(M, N, BS_R, BS_C, density, dtype="float32"):
Review comment:
Done.
--
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:
users@infra.apache.org
jcf94 commented on pull request #8065: URL: https://github.com/apache/tvm/pull/8065#issuecomment-842823573 Also cc @tkonolige -- 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: users@infra.apache.org
jcf94 commented on a change in pull request #8065:
URL: https://github.com/apache/tvm/pull/8065#discussion_r634068911
##########
File path: python/tvm/topi/sparse/utils.py
##########
@@ -116,11 +170,187 @@ def convert_model_dense_to_sparse(mod, params, random_params=False, bs_r=1, bs_c
Dict[Srting, tvm.nd.array]
The updated parameters.
"""
+
mod, params = ddo.simplify_fc_transpose.convert(mod["main"], params)
if random_params:
# Manually replace the parameters of dense model to sparse tensors
params = random_sparse_dense_params(mod, params, bs_r=bs_r, bs_c=bs_c, density=1 - sparsity)
+ params = random_sparse_conv2d_params(
+ mod, params, bs_r=bs_r, bs_c=bs_c, density=1 - sparsity, layout=layout
+ )
# Currently we only support to conver dense matmul to sparse dense matmul
Review comment:
Delete this line or add description about Conv2d.
--
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:
users@infra.apache.org
yuchaoli commented on a change in pull request #8065:
URL: https://github.com/apache/tvm/pull/8065#discussion_r634060352
##########
File path: tutorials/auto_scheduler/tune_network_x86.py
##########
@@ -138,7 +140,9 @@ def get_network(name, batch_size, layout="NHWC", dtype="float32", use_sparse=Fal
if use_sparse:
from tvm.topi.sparse.utils import convert_model_dense_to_sparse
- mod, params = convert_model_dense_to_sparse(mod, params, random_params=True)
+ mod, params = convert_model_dense_to_sparse(
+ mod, params, bs_r=sparse_block_size[0], bs_c=sparse_block_size[1], random_params=True
Review comment:
Done.
##########
File path: tutorials/auto_scheduler/tune_network_x86.py
##########
@@ -205,7 +215,33 @@ def run_tuning():
measure_callbacks=[auto_scheduler.RecordToFile(log_file)],
)
- tuner.tune(tune_option)
+ if use_sparse:
+ from tvm.topi.sparse.utils import (
+ sparse_conv2d_meet_condition_func,
+ sparse_conv2d_apply_func,
+ sparse_dense_meet_condition_func,
+ sparse_dense_apply_func,
+ )
+
+ search_policy = [
+ auto_scheduler.SketchPolicy(
+ task,
+ program_cost_model=auto_scheduler.XGBModel(),
+ init_search_callbacks=[
+ auto_scheduler.PreloadCustomSketchRule(
+ sparse_conv2d_meet_condition_func, sparse_conv2d_apply_func, "SparseConv2D"
+ ),
+ auto_scheduler.PreloadCustomSketchRule(
+ sparse_dense_meet_condition_func, sparse_dense_apply_func, "SparseDense"
+ ),
+ ],
+ )
+ for task in tasks
+ ]
Review comment:
Done.
--
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:
users@infra.apache.org
jcf94 merged pull request #8065: URL: https://github.com/apache/tvm/pull/8065 -- 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: users@infra.apache.org
jcf94 commented on a change in pull request #8065:
URL: https://github.com/apache/tvm/pull/8065#discussion_r634015717
##########
File path: tests/python/relay/test_sparse_conv2d_convert.py
##########
@@ -0,0 +1,110 @@
+# 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 itertools
+
+import numpy as np
+import scipy.sparse as sp
+
+
+import tvm
+from tvm.ir import IRModule
+from tvm import relay
+
+
+def random_bsr_matrix(M, N, BS_R, BS_C, density, dtype="float32"):
Review comment:
Use the `random_bsr_matrix` function in `topi/sparse/utils`, and we can also remove the the duplicate one in `tests/python/topi/python/test_topi_sparse.py`.
--
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:
users@infra.apache.org
tkonolige commented on a change in pull request #8065:
URL: https://github.com/apache/tvm/pull/8065#discussion_r634603256
##########
File path: python/tvm/relay/analysis/sparse_conv2d.py
##########
@@ -0,0 +1,154 @@
+# 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.
+# pylint: disable=no-else-return
+# pylint: disable=unidiomatic-typecheck
+"""
+This file contains helper functions for convert dense model
+to block sparse model
+"""
+from collections import namedtuple
+import numpy as np
+import scipy.sparse as sp
+import tvm
+from . import _ffi_api
+
+
+SparseAnalysisResult = namedtuple(
+ "SparseAnalysisResult",
+ [
+ "weight_name",
+ "weight_shape",
+ ],
+)
+
+
+def _search_conv2d_op_weight(expr):
+ """Search name of weight in all ```nn.conv2d``` operator
+ This is a helpful function to determine which param need
+ to be converted to sparse
+
+ Parameters
+ ----------
+ expr : relay.Expr
+ Expr will be searched
+
+ Returns
+ -------
+ ret : Array[String]
+ name of weight in all ``nn.conv2d``` operator
+ """
+ return _ffi_api.search_conv2d_op_weight(expr)
+
+
+def process_params(expr, params, block_size, sparsity_threshold, layout):
+ """[summary]
Review comment:
Missing documentation?
##########
File path: python/tvm/topi/nn/sparse.py
##########
@@ -478,6 +478,270 @@ def _traverse(t):
return sparse_input_map
+def _sparse_conv2d_bsr_compute_nhwc(data, weight_data, weight_indices, weight_indptr):
+ (m, h, w, k) = get_const_tuple(data.shape) # pylint: disable=C0103
+ if len(weight_data.shape) == 2:
+ _, bs_r = get_const_tuple(weight_data.shape)
+ elif len(weight_data.shape) == 3:
+ _, bs_r, bs_c = get_const_tuple(weight_data.shape)
+ (num_blocks_plus_1,) = get_const_tuple(weight_indptr.shape)
+ num_blocks = num_blocks_plus_1 - 1
+
+ def _compute_block(i, h, w, nb_j, j): # pylint: disable=C0103
+ row_start = weight_indptr[nb_j]
+ row_end = weight_indptr[nb_j + 1]
+ row_elems = row_end - row_start
+ elem_idx = te.reduce_axis((0, row_elems), name="elem_idx")
+ block_offset = row_start + elem_idx
+ block_j = weight_indices[block_offset]
+ if len(weight_data.shape) == 3:
+ c = te.reduce_axis((0, bs_c), name="c")
+ block_ij_val = weight_data[block_offset][j][c]
+ x_val = data[i, h, w, bs_c * block_j + c]
+ return te.sum(block_ij_val * x_val, axis=[elem_idx, c])
+ else:
+ block_ij_val = weight_data[block_offset][j]
+ x_val = data[i, h, w, block_j]
+ return te.sum(block_ij_val * x_val, axis=[elem_idx])
+
+ idxd = tvm.tir.indexdiv
+ idxm = tvm.tir.indexmod
+
+ bsrmm_block = te.compute(
+ (m, h, w, num_blocks, bs_r),
+ _compute_block,
+ tag="sparse_conv2d_sp_bsrmm_block",
+ attrs={"FLOP": 2 * m * num_blocks * bs_r * k * h * w},
+ )
+ return te.compute(
+ (m, h, w, num_blocks * bs_r),
+ lambda m, h, w, n: bsrmm_block[m, h, w, idxd(n, bs_r), idxm(n, bs_r)],
+ tag="sparse_conv2d_sp_bsrmm",
+ name="sparse_conv2d",
+ attrs={"layout": "NHWC"},
+ )
+
+
+def _sparse_conv2d_bsr_compute_nchw(data, weight_data, weight_indices, weight_indptr):
+ (m, k, h, w) = get_const_tuple(data.shape) # pylint: disable=C0103
+ if len(weight_data.shape) == 2:
+ _, bs_r = get_const_tuple(weight_data.shape)
+ elif len(weight_data.shape) == 3:
+ _, bs_r, bs_c = get_const_tuple(weight_data.shape)
+ (num_blocks_plus_1,) = get_const_tuple(weight_indptr.shape)
+ num_blocks = num_blocks_plus_1 - 1
+
+ def _compute_block(i, nb_j, j, h, w): # pylint: disable=C0103
+ row_start = weight_indptr[nb_j]
+ row_end = weight_indptr[nb_j + 1]
+ row_elems = row_end - row_start
+ elem_idx = te.reduce_axis((0, row_elems), name="elem_idx")
+ block_offset = row_start + elem_idx
+ block_j = weight_indices[block_offset]
+ if len(weight_data.shape) == 3:
+ c = te.reduce_axis((0, bs_c), name="c")
+ block_ij_val = weight_data[block_offset][j][c]
+ x_val = data[i, bs_c * block_j + c, h, w]
+ return te.sum(block_ij_val * x_val, axis=[elem_idx, c])
+ else:
+ block_ij_val = weight_data[block_offset][j]
+ x_val = data[i, block_j, h, w]
+ return te.sum(block_ij_val * x_val, axis=[elem_idx])
+
+ idxd = tvm.tir.indexdiv
+ idxm = tvm.tir.indexmod
+
+ bsrmm_block = te.compute(
+ (m, num_blocks, bs_r, h, w),
+ _compute_block,
+ tag="sparse_conv2d_sp_bsrmm_block",
+ attrs={"FLOP": 2 * m * num_blocks * bs_r * k * h * w},
+ )
+ return te.compute(
+ (m, num_blocks * bs_r, h, w),
+ lambda m, n, h, w: bsrmm_block[m, idxd(n, bs_r), idxm(n, bs_r), h, w],
+ tag="sparse_conv2d_sp_bsrmm",
+ name="sparse_conv2d",
+ attrs={"layout": "NCHW"},
+ )
+
+
+def sparse_conv2d(dense_data, sparse_data, sparse_indices, sparse_indptr, layout="NHWC"):
+ """
+ Computes sparse-conv2d(1*1) of `data` and
+ `(weight_data, weight_indices, weight_indptr)
+
+ Parameters
+ ----------
+ dense_data : tvm.te.Tensor
+ 4-D with shape [M, H, W, K], float32 (layout=NHWC)
+ 4-D with shape [M, K, H, W], float32 (layout=NCHW)
Review comment:
Is float32 required?
##########
File path: src/relay/transforms/convert_sparse_conv2d.cc
##########
@@ -0,0 +1,161 @@
+/*
+ * 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.
+ */
+
+/*!
+ *
+ * \file convert_sparse_conv2d.cc
+ *
+ * \brief Mutate dense operator to sparse dense operator
Review comment:
```suggestion
* \brief Mutate dense conv2d operator to sparse conv2d operator
```
##########
File path: src/relay/op/nn/sparse.cc
##########
@@ -237,5 +238,67 @@ RELAY_REGISTER_OP("nn.sparse_add")
.set_support_level(1)
.add_type_rel("SparseAdd", SparseAddRel);
+TVM_REGISTER_NODE_TYPE(SparseConv2DAttrs);
+
+bool SparseConv2dRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ ICHECK_EQ(types.size(), 5);
+ const auto* param = attrs.as<SparseConv2DAttrs>();
+ ICHECK(param != nullptr);
+
+ const auto* data = types[0].as<TensorTypeNode>();
+ const auto* weight_data = types[1].as<TensorTypeNode>();
+ ICHECK(weight_data->shape.size() == 1 || weight_data->shape.size() == 2 ||
+ weight_data->shape.size() == 3);
+ const auto* weight_indptr = types[3].as<TensorTypeNode>();
+ if (data == nullptr) return false;
+
+ if (weight_data->shape.size() == 2 || weight_data->shape.size() == 3) {
+ // BSR case.
+ if (param->layout == "NHWC") {
+ Array<IndexExpr> oshape({data->shape[0], data->shape[1], data->shape[2],
+ (weight_indptr->shape[0] - 1) * weight_data->shape[1]});
+ reporter->Assign(types[4], TensorType(oshape, data->dtype));
+ return true;
+ } else if (param->layout == "NCHW") {
+ Array<IndexExpr> oshape({data->shape[0],
+ (weight_indptr->shape[0] - 1) * weight_data->shape[1],
+ data->shape[2], data->shape[3]});
+ reporter->Assign(types[4], TensorType(oshape, data->dtype));
+ return true;
+ }
+ }
+ LOG(FATAL) << "Unknown weight ndim for nn.sparse_conv2d, should be 3 (BSR)";
Review comment:
Can you print out the actual weight here.
##########
File path: src/relay/op/nn/sparse.cc
##########
@@ -237,5 +238,67 @@ RELAY_REGISTER_OP("nn.sparse_add")
.set_support_level(1)
.add_type_rel("SparseAdd", SparseAddRel);
+TVM_REGISTER_NODE_TYPE(SparseConv2DAttrs);
+
+bool SparseConv2dRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ ICHECK_EQ(types.size(), 5);
Review comment:
Can you add an error message.
--
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:
users@infra.apache.org
jcf94 commented on a change in pull request #8065:
URL: https://github.com/apache/tvm/pull/8065#discussion_r634013363
##########
File path: tutorials/auto_scheduler/tune_network_x86.py
##########
@@ -138,7 +140,9 @@ def get_network(name, batch_size, layout="NHWC", dtype="float32", use_sparse=Fal
if use_sparse:
from tvm.topi.sparse.utils import convert_model_dense_to_sparse
- mod, params = convert_model_dense_to_sparse(mod, params, random_params=True)
+ mod, params = convert_model_dense_to_sparse(
+ mod, params, bs_r=sparse_block_size[0], bs_c=sparse_block_size[1], random_params=True
Review comment:
Let's keep this tutorial simple.
Suggest to move the sparse_block_size setting inside this function, for example directly set the block size here.
So that we can avoid to modify the function parameters and add `sparse_block_size = (8, 1)` outside.
##########
File path: tutorials/auto_scheduler/tune_network_x86.py
##########
@@ -205,7 +215,33 @@ def run_tuning():
measure_callbacks=[auto_scheduler.RecordToFile(log_file)],
)
- tuner.tune(tune_option)
+ if use_sparse:
+ from tvm.topi.sparse.utils import (
+ sparse_conv2d_meet_condition_func,
+ sparse_conv2d_apply_func,
+ sparse_dense_meet_condition_func,
+ sparse_dense_apply_func,
+ )
+
+ search_policy = [
+ auto_scheduler.SketchPolicy(
+ task,
+ program_cost_model=auto_scheduler.XGBModel(),
+ init_search_callbacks=[
+ auto_scheduler.PreloadCustomSketchRule(
+ sparse_conv2d_meet_condition_func, sparse_conv2d_apply_func, "SparseConv2D"
+ ),
+ auto_scheduler.PreloadCustomSketchRule(
+ sparse_dense_meet_condition_func, sparse_dense_apply_func, "SparseDense"
+ ),
+ ],
+ )
+ for task in tasks
+ ]
Review comment:
How about to pack these 4 condition func & apply func to a single entrance. For example:
```python
def sparse_sketch_rules() -> List[PreloadCustomSketchRule]
...
...
init_search_callbacks=sparse_sketch_rules(),
...
```
--
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:
users@infra.apache.org
yuchaoli commented on a change in pull request #8065:
URL: https://github.com/apache/tvm/pull/8065#discussion_r634076852
##########
File path: python/tvm/topi/sparse/utils.py
##########
@@ -116,11 +170,187 @@ def convert_model_dense_to_sparse(mod, params, random_params=False, bs_r=1, bs_c
Dict[Srting, tvm.nd.array]
The updated parameters.
"""
+
mod, params = ddo.simplify_fc_transpose.convert(mod["main"], params)
if random_params:
# Manually replace the parameters of dense model to sparse tensors
params = random_sparse_dense_params(mod, params, bs_r=bs_r, bs_c=bs_c, density=1 - sparsity)
+ params = random_sparse_conv2d_params(
+ mod, params, bs_r=bs_r, bs_c=bs_c, density=1 - sparsity, layout=layout
+ )
# Currently we only support to conver dense matmul to sparse dense matmul
Review comment:
Done.
--
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:
users@infra.apache.org
jcf94 commented on pull request #8065: URL: https://github.com/apache/tvm/pull/8065#issuecomment-843986728 Merged. Thanks! @yuchaoli @tkonolige -- 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: users@infra.apache.org