From 05f9e36cc476244beeef3512f2e3d86d35766327 Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Wed, 10 Feb 2021 21:08:20 +0530 Subject: [PATCH 1/9] [TOPI] Sparse Add Op added --- python/tvm/relay/op/strategy/generic.py | 28 +++++++++ python/tvm/topi/generic/nn.py | 15 +++++ python/tvm/topi/nn/sparse.py | 65 ++++++++++++++++++++ src/relay/op/nn/sparse.cc | 41 ++++++++++++ tests/python/topi/python/test_topi_sparse.py | 24 ++++++++ 5 files changed, 173 insertions(+) diff --git a/python/tvm/relay/op/strategy/generic.py b/python/tvm/relay/op/strategy/generic.py index be86ea9d9184..426923895bce 100644 --- a/python/tvm/relay/op/strategy/generic.py +++ b/python/tvm/relay/op/strategy/generic.py @@ -799,6 +799,34 @@ def sparse_dense_padded_strategy(attrs, inputs, out_type, target): raise NotImplementedError("sparse_dense_padded is only implemented for cuda") +# sparse_add +def wrap_compute_sparse_add(topi_compute): + """wrap sparse add topi compute""" + + def _compute_sparse_add(attrs, inputs, out_type): + return [topi_compute(inputs[0], inputs[1], inputs[2], inputs[3])] + + return _compute_sparse_add + + +@override_native_generic_func("sparse_add_strategy") +def sparse_add_strategy(attrs, inputs, out_type, target): + """sparse add generic strategy""" + logger.warning("sparse add is not optimized for this platform.") + strategy = _op.OpStrategy() + strategy.add_implementation( + wrap_compute_sparse_add(topi.nn.sparse_add), + wrap_topi_schedule(topi.generic.schedule_sparse_add), + name="sparse_add.generic", + ) + return strategy + +@generic_func +def schedule_sparse_add(attrs, outs, target): + """schedule sparse_add""" + with target: + return topi.generic.schedule_sparse_add(outs) + # sparse_transpose @generic_func def schedule_sparse_transpose(attrs, outs, target): diff --git a/python/tvm/topi/generic/nn.py b/python/tvm/topi/generic/nn.py index 60ccd0d36abf..cbb6a94a2819 100644 --- a/python/tvm/topi/generic/nn.py +++ b/python/tvm/topi/generic/nn.py @@ -729,6 +729,21 @@ def schedule_sparse_transpose(outs): """ return _default_schedule(outs, False) +def schedule_sparse_add(outs): + """Schedule for sparse_add + + Parameters + ---------- + outs: Array of Tensor + The computation graph description of sparse_add + in the format of an array of tensors. + + Returns + ------- + sch: Schedule + The computation schedule for the op. + """ + return _default_schedule(outs, False) def schedule_batch_matmul(outs): """Schedule for batch_matmul diff --git a/python/tvm/topi/nn/sparse.py b/python/tvm/topi/nn/sparse.py index 1bf18df09da3..293a414098e1 100644 --- a/python/tvm/topi/nn/sparse.py +++ b/python/tvm/topi/nn/sparse.py @@ -468,3 +468,68 @@ def _traverse(t): sparse_input_map[sparse_indptr] = sparse_prefix + "W_indptr" return sparse_input_map + + +def sparse_add(dense_data, sparse_data, sparse_indices, sparse_indptr): + """ + Computes sparse-dense addition + + Parameters + ---------- + dense_data : tvm.te.Tensor + 2-D with shape [M, K], float32 + + sparse_data : tvm.te.Tensor + 1-D with shape [nnz] (CSR) or + + sparse_indices : tvm.te.Tensor + 1-D with shape [nnz] (CSR) or + + sparse_indptr : tvm.te.Tensor + 1-D with shape [N + 1] (CSR) or + + Returns + ------- + output : tvm.te.Tensor + 2-D with shape [M, N] + """ + #TODO(ANSHUMAN87): support BSR format too + assert len(sparse_data.shape) == 1 + return _sparse_add_csr(dense_data, sparse_data, sparse_indices, sparse_indptr) + +def _sparse_add_csr(dense_data_inp, sparse_data_inp, sparse_indices_inp, sparse_indptr_inp): + oshape = get_const_tuple(dense_data_inp.shape) + + def _csr_add_ir(dense_data, sparse_data, sparse_indices, sparse_indptr, out_data): + irb = tvm.tir.ir_builder.create() + dense_data_ptr = irb.buffer_ptr(dense_data) + sparse_data_ptr = irb.buffer_ptr(sparse_data) + sparse_indices_ptr = irb.buffer_ptr(sparse_indices) + sparse_indptr_ptr = irb.buffer_ptr(sparse_indptr) + + out_data_ptr = irb.buffer_ptr(out_data) + + with irb.for_range(0, oshape[0], kind="serial", name="row") as row: + with irb.for_range(0, oshape[1], kind="serial", name="col") as col: + out_data_ptr[row, col] = dense_data_ptr[row, col] + + with irb.for_range(0, oshape[0], kind="serial", name="row") as row: + offset = sparse_indptr_ptr[row] + diff = sparse_indptr_ptr[row + 1] - sparse_indptr_ptr[row] + with irb.for_range(0, diff, kind="serial", name="idx") as idx: + real_idx = offset + idx + col = sparse_indices_ptr[real_idx] + out_data_ptr[row, col] = sparse_data_ptr[real_idx] + out_data_ptr[row, col] + + return irb.get() + + return te.extern( + shape=oshape, + inputs=[dense_data_inp, sparse_data_inp, sparse_indices_inp, sparse_indptr_inp], + fcompute=lambda ins, outs: _csr_add_ir( + ins[0], ins[1], ins[2], ins[3], outs[0] + ), + tag="sparse_add_csr", + dtype=["float32", "float32", "int32", "int32"], + name="out", + ) diff --git a/src/relay/op/nn/sparse.cc b/src/relay/op/nn/sparse.cc index 6322cfffd7c2..7b3aaa285b03 100644 --- a/src/relay/op/nn/sparse.cc +++ b/src/relay/op/nn/sparse.cc @@ -196,5 +196,46 @@ RELAY_REGISTER_OP("nn.sparse_transpose") .set_support_level(1) .add_type_rel("SparseTranspose", SparseTransposeRel); +// relay.nn.sparse_add +bool SparseAddRel(const Array& types, int num_inputs, const Attrs& attrs, + const TypeReporter& reporter) { + ICHECK_EQ(types.size(), 5); + const auto* dense_data = types[0].as(); + const auto* sparse_data = types[1].as(); + ICHECK_EQ(sparse_data->shape.size(), 1); + const auto* sparse_indices = types[2].as(); + ICHECK_EQ(sparse_indices->shape.size(), 1); + const auto* sparse_indptr = types[3].as(); + + reporter->Assign(types[4], TensorType(dense_data->shape, dense_data->dtype)); + return true; +} + +Expr MakeSparseAdd(Expr dense_data, Expr sparse_data, Expr sparse_indices, Expr sparse_indptr) { + static const Op& op = Op::Get("nn.sparse_add"); + return Call(op, {dense_data, sparse_data, sparse_indices, sparse_indptr}, Attrs(), {}); +} + +TVM_REGISTER_GLOBAL("relay.op.nn._make.sparse_add").set_body_typed(MakeSparseAdd); + +RELAY_REGISTER_OP("nn.sparse_add") + .describe(R"code(Add a dense matrix X with sparse matrix Y. Only support square sparse matrix + +- **dense**: `(N, N)` +- **sparse**: `(N, N)` + +- **out**: `(N, N)`. + +)code" TVM_ADD_FILELINE) + .set_attrs_type() + .set_num_inputs(4) + .add_argument("dense_data", "2D Tensor", "Dense data matrix.") + .add_argument("sparse_data", "1D Tensor", "Sparse data matrix.") + .add_argument("sparse_indices", "1D Tensor", "Sparse indices matrix.") + .add_argument("sparse_indptr", "1D Tensor", "Sparse index pointer matrix.") + .set_support_level(1) + .add_type_rel("SparseAdd", SparseAddRel); + + } // namespace relay } // namespace tvm diff --git a/tests/python/topi/python/test_topi_sparse.py b/tests/python/topi/python/test_topi_sparse.py index d5bd7aa1a21e..1b3bcc7f2732 100644 --- a/tests/python/topi/python/test_topi_sparse.py +++ b/tests/python/topi/python/test_topi_sparse.py @@ -525,6 +525,29 @@ def test_sparse_dense_padded_alter_op(): with tvm.transform.PassContext(opt_level=3, required_pass="AlterOpLayout"): x = relay.build(tvm.IRModule.from_expr(f), target=tvm.target.Target("cuda")) +def test_sparse_add_csr(): + M, K, density = 3, 49, 0.2 + X_np = np.random.randn(M, K).astype("float32") + Y_sp_np = sp.random(M, K, density=density, format="csr", dtype="float32") + Y_np = Y_sp_np.todense() + Z_np = X_np + Y_np + + Y_data = te.placeholder(shape=Y_sp_np.data.shape, dtype=str(Y_sp_np.data.dtype)) + Y_indices = te.placeholder(shape=Y_sp_np.indices.shape, dtype=str(Y_sp_np.indices.dtype)) + Y_indptr = te.placeholder(shape=Y_sp_np.indptr.shape, dtype=str(Y_sp_np.indptr.dtype)) + X = te.placeholder(shape=X_np.shape, dtype=str(X_np.dtype)) + Z = topi.nn.sparse_add(X, Y_data, Y_indices, Y_indptr) + s = te.create_schedule(Z.op) + func = tvm.build(s, [X, Y_data, Y_indices, Y_indptr, Z]) + Z_tvm = tvm.nd.array(np.zeros(Z_np.shape, dtype=Z_np.dtype)) + func( + tvm.nd.array(X_np), + tvm.nd.array(Y_sp_np.data), + tvm.nd.array(Y_sp_np.indices), + tvm.nd.array(Y_sp_np.indptr), + Z_tvm, + ) + tvm.testing.assert_allclose(Z_tvm.asnumpy(), Z_np, atol=1e-4, rtol=1e-4) if __name__ == "__main__": test_csrmv() @@ -537,3 +560,4 @@ def test_sparse_dense_padded_alter_op(): test_sparse_dense_padded_alter_op() test_sparse_dense_csr_reverse() test_sparse_dense_bsr_reverse() + test_sparse_add_csr() From 8be54938abb3c26f55ec2b8fce7e2b8ee536e5cc Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Sat, 13 Feb 2021 10:33:36 +0530 Subject: [PATCH 2/9] lint resolved --- python/tvm/relay/op/strategy/generic.py | 2 ++ python/tvm/topi/generic/nn.py | 2 ++ python/tvm/topi/nn/sparse.py | 7 +++---- src/relay/op/nn/sparse.cc | 3 +-- tests/python/topi/python/test_topi_sparse.py | 2 ++ 5 files changed, 10 insertions(+), 6 deletions(-) diff --git a/python/tvm/relay/op/strategy/generic.py b/python/tvm/relay/op/strategy/generic.py index 426923895bce..ff9238e29485 100644 --- a/python/tvm/relay/op/strategy/generic.py +++ b/python/tvm/relay/op/strategy/generic.py @@ -821,12 +821,14 @@ def sparse_add_strategy(attrs, inputs, out_type, target): ) return strategy + @generic_func def schedule_sparse_add(attrs, outs, target): """schedule sparse_add""" with target: return topi.generic.schedule_sparse_add(outs) + # sparse_transpose @generic_func def schedule_sparse_transpose(attrs, outs, target): diff --git a/python/tvm/topi/generic/nn.py b/python/tvm/topi/generic/nn.py index cbb6a94a2819..49281e0356a6 100644 --- a/python/tvm/topi/generic/nn.py +++ b/python/tvm/topi/generic/nn.py @@ -729,6 +729,7 @@ def schedule_sparse_transpose(outs): """ return _default_schedule(outs, False) + def schedule_sparse_add(outs): """Schedule for sparse_add @@ -745,6 +746,7 @@ def schedule_sparse_add(outs): """ return _default_schedule(outs, False) + def schedule_batch_matmul(outs): """Schedule for batch_matmul diff --git a/python/tvm/topi/nn/sparse.py b/python/tvm/topi/nn/sparse.py index 293a414098e1..dd43113172cc 100644 --- a/python/tvm/topi/nn/sparse.py +++ b/python/tvm/topi/nn/sparse.py @@ -493,10 +493,11 @@ def sparse_add(dense_data, sparse_data, sparse_indices, sparse_indptr): output : tvm.te.Tensor 2-D with shape [M, N] """ - #TODO(ANSHUMAN87): support BSR format too + # TODO(ANSHUMAN87): support BSR format too assert len(sparse_data.shape) == 1 return _sparse_add_csr(dense_data, sparse_data, sparse_indices, sparse_indptr) + def _sparse_add_csr(dense_data_inp, sparse_data_inp, sparse_indices_inp, sparse_indptr_inp): oshape = get_const_tuple(dense_data_inp.shape) @@ -526,9 +527,7 @@ def _csr_add_ir(dense_data, sparse_data, sparse_indices, sparse_indptr, out_data return te.extern( shape=oshape, inputs=[dense_data_inp, sparse_data_inp, sparse_indices_inp, sparse_indptr_inp], - fcompute=lambda ins, outs: _csr_add_ir( - ins[0], ins[1], ins[2], ins[3], outs[0] - ), + fcompute=lambda ins, outs: _csr_add_ir(ins[0], ins[1], ins[2], ins[3], outs[0]), tag="sparse_add_csr", dtype=["float32", "float32", "int32", "int32"], name="out", diff --git a/src/relay/op/nn/sparse.cc b/src/relay/op/nn/sparse.cc index 7b3aaa285b03..f026aa949bd9 100644 --- a/src/relay/op/nn/sparse.cc +++ b/src/relay/op/nn/sparse.cc @@ -198,7 +198,7 @@ RELAY_REGISTER_OP("nn.sparse_transpose") // relay.nn.sparse_add bool SparseAddRel(const Array& types, int num_inputs, const Attrs& attrs, - const TypeReporter& reporter) { + const TypeReporter& reporter) { ICHECK_EQ(types.size(), 5); const auto* dense_data = types[0].as(); const auto* sparse_data = types[1].as(); @@ -236,6 +236,5 @@ RELAY_REGISTER_OP("nn.sparse_add") .set_support_level(1) .add_type_rel("SparseAdd", SparseAddRel); - } // namespace relay } // namespace tvm diff --git a/tests/python/topi/python/test_topi_sparse.py b/tests/python/topi/python/test_topi_sparse.py index 1b3bcc7f2732..5d92694cfaa4 100644 --- a/tests/python/topi/python/test_topi_sparse.py +++ b/tests/python/topi/python/test_topi_sparse.py @@ -525,6 +525,7 @@ def test_sparse_dense_padded_alter_op(): with tvm.transform.PassContext(opt_level=3, required_pass="AlterOpLayout"): x = relay.build(tvm.IRModule.from_expr(f), target=tvm.target.Target("cuda")) + def test_sparse_add_csr(): M, K, density = 3, 49, 0.2 X_np = np.random.randn(M, K).astype("float32") @@ -549,6 +550,7 @@ def test_sparse_add_csr(): ) tvm.testing.assert_allclose(Z_tvm.asnumpy(), Z_np, atol=1e-4, rtol=1e-4) + if __name__ == "__main__": test_csrmv() test_csrmm() From 94791dd9c2e65407e4d30330fd4b520a362d75a7 Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Sat, 13 Feb 2021 12:59:36 +0530 Subject: [PATCH 3/9] TF frontend support added --- python/tvm/relay/frontend/tensorflow.py | 33 ++++++++++++++ python/tvm/relay/op/nn/_nn.py | 5 +++ python/tvm/relay/op/nn/nn.py | 33 ++++++++++++++ python/tvm/topi/nn/sparse.py | 4 +- src/relay/op/nn/sparse.cc | 10 ++--- .../frontend/tensorflow/test_forward.py | 44 +++++++++++++++++++ 6 files changed, 121 insertions(+), 8 deletions(-) diff --git a/python/tvm/relay/frontend/tensorflow.py b/python/tvm/relay/frontend/tensorflow.py index c79c495b0360..cbc8a10ef8fa 100644 --- a/python/tvm/relay/frontend/tensorflow.py +++ b/python/tvm/relay/frontend/tensorflow.py @@ -1286,6 +1286,38 @@ def _impl(inputs, attr, params, mod): return _impl +def _sparse_tensor_dense_add(): + # Sparse utility from scipy + from scipy.sparse import csr_matrix + + def _impl(inputs, attr, params, mod): + assert len(inputs) == 4, "There should be 4 input tensors" + + indices_tensor = _infer_value(inputs[0], params, mod).asnumpy() + values_tensor = _infer_value(inputs[1], params, mod).asnumpy() + dense_shape_tensor = _infer_value(inputs[2], params, mod).asnumpy() + + data = inputs[3] + + rows = [x[0] for x in indices_tensor] + cols = [x[1] for x in indices_tensor] + + # Create scipy sparse Tensor(CSR) + weight_sp = csr_matrix( + (values_tensor, (rows, cols)), shape=tuple(dense_shape_tensor.tolist()) + ) + + weight_data = _expr.const(weight_sp.data, weight_sp.data.dtype) + weight_indptrs = _expr.const(weight_sp.indptr, weight_sp.indptr.dtype) + weight_indices = _expr.const(weight_sp.indices, weight_sp.indices.dtype) + + ret = _op.nn.sparse_add(data, [weight_data, weight_indices, weight_indptrs]) + + return ret + + return _impl + + def _identity(): def _impl(inputs, attr, params, mod): return inputs[0] @@ -2787,6 +2819,7 @@ def _impl(inputs, attr, params, mod): "SparseSegmentSqrtNWithNumSegments": _sparse_segment_sum_sqrtn_with_num_segments(), "SparseSegmentMean": _sparse_segment_mean(), "SparseSegmentMeanWithNumSegments": _sparse_segment_mean_with_num_segments(), + "SparseTensorDenseAdd": _sparse_tensor_dense_add(), "Split": _split(False), "SplitV": _split(True), "Sqrt": AttrCvt("sqrt"), diff --git a/python/tvm/relay/op/nn/_nn.py b/python/tvm/relay/op/nn/_nn.py index 6ae86c0786e5..e4cceecc1890 100644 --- a/python/tvm/relay/op/nn/_nn.py +++ b/python/tvm/relay/op/nn/_nn.py @@ -142,6 +142,11 @@ def alter_op_layout_sparse_dense(attrs, inputs, tinfos, out_type): return topi.nn.sparse_dense_alter_layout(attrs, inputs, tinfos, out_type) +# sparse_add +reg.register_strategy("nn.sparse_add", strategy.sparse_add_strategy) +reg.register_pattern("nn.sparse_add", reg.OpPattern.ELEMWISE) + + @reg.register_compute("nn.internal.sparse_dense_padded") def compute_sparse_dense_padded(attrs, inputs, out_type): """Compute definition of sparse_dense_padded""" diff --git a/python/tvm/relay/op/nn/nn.py b/python/tvm/relay/op/nn/nn.py index 5135ac74de25..2b854c840dc4 100644 --- a/python/tvm/relay/op/nn/nn.py +++ b/python/tvm/relay/op/nn/nn.py @@ -2148,6 +2148,39 @@ def sparse_transpose(x): return expr.TupleWrapper(_make.sparse_transpose(x[0], x[1], x[2]), 3) +# pylint: disable=no-else-return,inconsistent-return-statements +def sparse_add(dense_mat, sparse_mat): + r""" + Computes the matrix addition of `dense_mat` and `sparse_mat`, where `dense_mat` is + a dense matrix and `sparse_mat` is a sparse (either BSR or CSR) namedtuple with + fields `data`, `indices`, and `indptr`. + + .. math:: + + \mbox{sparse_add}(dense_mat, sparse_mat)[m, n] = \mbox{add}(\mbox{as_dense}(S), (D))[m, n] + + where `as_dense` returns dense equivalent of the given S(sparse matrix) + while performing addition with given D(dense matrix). + + Parameters + ---------- + dense_mat : tvm.relay.Expr + The input dense matrix for the matrix multiplication + + sparse_mat : Union[namedtuple, Tuple[ndarray, ndarray, ndarray]]. + The input sparse matrix for the matrix multiplication. + + Returns + ------- + result: tvm.relay.Expr + The computed result. + """ + if hasattr(sparse_mat, "indices"): + return _make.sparse_add(dense_mat, sparse_mat.data, sparse_mat.indices, sparse_mat.indptr) + else: + return _make.sparse_add(dense_mat, sparse_mat[0], sparse_mat[1], sparse_mat[2]) + + def contrib_conv2d_winograd_without_weight_transform( data, weight, diff --git a/python/tvm/topi/nn/sparse.py b/python/tvm/topi/nn/sparse.py index dd43113172cc..3784aab5df4d 100644 --- a/python/tvm/topi/nn/sparse.py +++ b/python/tvm/topi/nn/sparse.py @@ -477,7 +477,7 @@ def sparse_add(dense_data, sparse_data, sparse_indices, sparse_indptr): Parameters ---------- dense_data : tvm.te.Tensor - 2-D with shape [M, K], float32 + 2-D with shape [M, N], float32 sparse_data : tvm.te.Tensor 1-D with shape [nnz] (CSR) or @@ -486,7 +486,7 @@ def sparse_add(dense_data, sparse_data, sparse_indices, sparse_indptr): 1-D with shape [nnz] (CSR) or sparse_indptr : tvm.te.Tensor - 1-D with shape [N + 1] (CSR) or + 1-D with shape [M + 1] (CSR) or Returns ------- diff --git a/src/relay/op/nn/sparse.cc b/src/relay/op/nn/sparse.cc index f026aa949bd9..08a8394858a1 100644 --- a/src/relay/op/nn/sparse.cc +++ b/src/relay/op/nn/sparse.cc @@ -205,7 +205,6 @@ bool SparseAddRel(const Array& types, int num_inputs, const Attrs& attrs, ICHECK_EQ(sparse_data->shape.size(), 1); const auto* sparse_indices = types[2].as(); ICHECK_EQ(sparse_indices->shape.size(), 1); - const auto* sparse_indptr = types[3].as(); reporter->Assign(types[4], TensorType(dense_data->shape, dense_data->dtype)); return true; @@ -219,15 +218,14 @@ Expr MakeSparseAdd(Expr dense_data, Expr sparse_data, Expr sparse_indices, Expr TVM_REGISTER_GLOBAL("relay.op.nn._make.sparse_add").set_body_typed(MakeSparseAdd); RELAY_REGISTER_OP("nn.sparse_add") - .describe(R"code(Add a dense matrix X with sparse matrix Y. Only support square sparse matrix + .describe(R"code(Add a dense matrix X with sparse matrix Y. -- **dense**: `(N, N)` -- **sparse**: `(N, N)` +- **dense**: `(M, N)` +- **sparse**: `(M, N)` -- **out**: `(N, N)`. +- **out**: `(M, N)`. )code" TVM_ADD_FILELINE) - .set_attrs_type() .set_num_inputs(4) .add_argument("dense_data", "2D Tensor", "Dense data matrix.") .add_argument("sparse_data", "1D Tensor", "Sparse data matrix.") diff --git a/tests/python/frontend/tensorflow/test_forward.py b/tests/python/frontend/tensorflow/test_forward.py index 81aeb5ef886c..9de88a721da4 100644 --- a/tests/python/frontend/tensorflow/test_forward.py +++ b/tests/python/frontend/tensorflow/test_forward.py @@ -2352,6 +2352,50 @@ def test_forward_sparse_to_dense_v2(): _test_sparse_to_dense_v2([[0, 0], [1, 3], [4, 3]], [3.0, 6.0, 9.0], [5, 5], "float32", 1.9) +####################################################################### +# tensorflow.sparse.add +# ---------------------------------- + + +def _test_sparse_add(indices, values, A_shape, B_shape, dtype, flip=False): + """ One iteration of tf.sparse.add """ + + # TODO(ANSHUMAN87): support cuda + # TODO(ANSHUMAN87): support flip case + # TODO(ANSHUMAN87): support both sparse input case + + with tf.Graph().as_default(): + A_sp = tf.sparse.SparseTensor(indices=indices, values=values, dense_shape=A_shape) + B = tf.placeholder(shape=B_shape, dtype=dtype, name="B") + + if flip: + result = tf.sparse.add(B, A_sp, threshold=0) + else: + result = tf.sparse.add(A_sp, B, threshold=0) + + B_np = np.random.uniform(high=5.0, size=B_shape).astype(dtype) + + compare_tf_with_tvm([B_np], [B.name], result.name, no_gpu=True) + + +def test_sparse_add(): + """ sparse.add op test""" + ################################################################### + # + # In order to create a SparseTensor, it requires 3 input as below: + # SparseTensor(indices=[[0, 0], [1, 2]], values=[1, 2], dense_shape=[3, 4]) + # + # Above Sparse can be represented in Dense as below : + # [[1, 0, 0, 0] + # [0, 0, 2, 0] + # [0, 0, 0, 0]] + # + # ------------------------------------------------------------------ + + # TODO(ANSHUMAN87): add more test case + _test_sparse_add([[0, 0], [1, 2]], [4.0, 8.0], [3, 4], [3, 4], "float32") + + ####################################################################### # StridedSlice # ------------ From 83fd6660478cbb4f107cafa99eb15120849b55b7 Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Wed, 17 Feb 2021 22:08:08 +0530 Subject: [PATCH 4/9] Test case added --- tests/python/frontend/tensorflow/test_forward.py | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/tests/python/frontend/tensorflow/test_forward.py b/tests/python/frontend/tensorflow/test_forward.py index 9de88a721da4..6b0b13a6bc73 100644 --- a/tests/python/frontend/tensorflow/test_forward.py +++ b/tests/python/frontend/tensorflow/test_forward.py @@ -2361,13 +2361,13 @@ def _test_sparse_add(indices, values, A_shape, B_shape, dtype, flip=False): """ One iteration of tf.sparse.add """ # TODO(ANSHUMAN87): support cuda - # TODO(ANSHUMAN87): support flip case # TODO(ANSHUMAN87): support both sparse input case with tf.Graph().as_default(): A_sp = tf.sparse.SparseTensor(indices=indices, values=values, dense_shape=A_shape) B = tf.placeholder(shape=B_shape, dtype=dtype, name="B") + # TODO(ANSHUMAN87): support user input threashold values if flip: result = tf.sparse.add(B, A_sp, threshold=0) else: @@ -2392,8 +2392,10 @@ def test_sparse_add(): # # ------------------------------------------------------------------ - # TODO(ANSHUMAN87): add more test case _test_sparse_add([[0, 0], [1, 2]], [4.0, 8.0], [3, 4], [3, 4], "float32") + _test_sparse_add([[0, 0], [1, 2]], [4.0, 8.0], [3, 4], [3, 4], "float32", True) + _test_sparse_add([[0, 0], [1, 3], [4, 3]], [3.0, 6.0, 9.0], [5, 5], [5, 5], "float32") + _test_sparse_add([[0, 0], [1, 3], [4, 3]], [3.0, 6.0, 9.0], [5, 5], [5, 5], "float32", True) ####################################################################### From 16fd09444d94d17ef4d8497cf9d16ae0fea07da3 Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Thu, 18 Feb 2021 22:42:53 +0530 Subject: [PATCH 5/9] [1] Review comment handled --- python/tvm/relay/frontend/tensorflow.py | 4 +++- python/tvm/relay/op/nn/_nn.py | 2 +- python/tvm/relay/op/strategy/generic.py | 2 +- python/tvm/topi/nn/sparse.py | 9 +++++++-- src/relay/op/nn/sparse.cc | 11 ++++++----- 5 files changed, 18 insertions(+), 10 deletions(-) diff --git a/python/tvm/relay/frontend/tensorflow.py b/python/tvm/relay/frontend/tensorflow.py index cbc8a10ef8fa..f56d187b6a63 100644 --- a/python/tvm/relay/frontend/tensorflow.py +++ b/python/tvm/relay/frontend/tensorflow.py @@ -1291,7 +1291,9 @@ def _sparse_tensor_dense_add(): from scipy.sparse import csr_matrix def _impl(inputs, attr, params, mod): - assert len(inputs) == 4, "There should be 4 input tensors" + assert ( + len(inputs) == 4 + ), "There should be 4 input tensors [sparse_indices, sparse_values, sparse_shape, dense]." indices_tensor = _infer_value(inputs[0], params, mod).asnumpy() values_tensor = _infer_value(inputs[1], params, mod).asnumpy() diff --git a/python/tvm/relay/op/nn/_nn.py b/python/tvm/relay/op/nn/_nn.py index e4cceecc1890..af64873ee904 100644 --- a/python/tvm/relay/op/nn/_nn.py +++ b/python/tvm/relay/op/nn/_nn.py @@ -144,7 +144,7 @@ def alter_op_layout_sparse_dense(attrs, inputs, tinfos, out_type): # sparse_add reg.register_strategy("nn.sparse_add", strategy.sparse_add_strategy) -reg.register_pattern("nn.sparse_add", reg.OpPattern.ELEMWISE) +reg.register_pattern("nn.sparse_add", reg.OpPattern.OPAQUE) @reg.register_compute("nn.internal.sparse_dense_padded") diff --git a/python/tvm/relay/op/strategy/generic.py b/python/tvm/relay/op/strategy/generic.py index ff9238e29485..f2775d07ea7c 100644 --- a/python/tvm/relay/op/strategy/generic.py +++ b/python/tvm/relay/op/strategy/generic.py @@ -816,7 +816,7 @@ def sparse_add_strategy(attrs, inputs, out_type, target): strategy = _op.OpStrategy() strategy.add_implementation( wrap_compute_sparse_add(topi.nn.sparse_add), - wrap_topi_schedule(topi.generic.schedule_sparse_add), + wrap_topi_schedule(topi.generic.schedule_extern), name="sparse_add.generic", ) return strategy diff --git a/python/tvm/topi/nn/sparse.py b/python/tvm/topi/nn/sparse.py index 3784aab5df4d..de46117fa811 100644 --- a/python/tvm/topi/nn/sparse.py +++ b/python/tvm/topi/nn/sparse.py @@ -494,7 +494,7 @@ def sparse_add(dense_data, sparse_data, sparse_indices, sparse_indptr): 2-D with shape [M, N] """ # TODO(ANSHUMAN87): support BSR format too - assert len(sparse_data.shape) == 1 + assert len(sparse_data.shape) == 1, "only CSR format is supported" return _sparse_add_csr(dense_data, sparse_data, sparse_indices, sparse_indptr) @@ -529,6 +529,11 @@ def _csr_add_ir(dense_data, sparse_data, sparse_indices, sparse_indptr, out_data inputs=[dense_data_inp, sparse_data_inp, sparse_indices_inp, sparse_indptr_inp], fcompute=lambda ins, outs: _csr_add_ir(ins[0], ins[1], ins[2], ins[3], outs[0]), tag="sparse_add_csr", - dtype=["float32", "float32", "int32", "int32"], + dtype=[ + dense_data_inp.dtype, + sparse_data_inp.dtype, + sparse_indices_inp.dtype, + sparse_indptr_inp.dtype, + ], name="out", ) diff --git a/src/relay/op/nn/sparse.cc b/src/relay/op/nn/sparse.cc index 08a8394858a1..94d088d08e61 100644 --- a/src/relay/op/nn/sparse.cc +++ b/src/relay/op/nn/sparse.cc @@ -202,9 +202,10 @@ bool SparseAddRel(const Array& types, int num_inputs, const Attrs& attrs, ICHECK_EQ(types.size(), 5); const auto* dense_data = types[0].as(); const auto* sparse_data = types[1].as(); - ICHECK_EQ(sparse_data->shape.size(), 1); + ICHECK(reporter->Assert(sparse_data->dtype == dense_data->dtype)); + ICHECK(reporter->Assert(sparse_data->shape.size() == 1)); const auto* sparse_indices = types[2].as(); - ICHECK_EQ(sparse_indices->shape.size(), 1); + ICHECK(reporter->Assert(sparse_indices->shape.size() == 1)); reporter->Assign(types[4], TensorType(dense_data->shape, dense_data->dtype)); return true; @@ -228,9 +229,9 @@ RELAY_REGISTER_OP("nn.sparse_add") )code" TVM_ADD_FILELINE) .set_num_inputs(4) .add_argument("dense_data", "2D Tensor", "Dense data matrix.") - .add_argument("sparse_data", "1D Tensor", "Sparse data matrix.") - .add_argument("sparse_indices", "1D Tensor", "Sparse indices matrix.") - .add_argument("sparse_indptr", "1D Tensor", "Sparse index pointer matrix.") + .add_argument("sparse_data", "1D Tensor", "Sparse data vector.") + .add_argument("sparse_indices", "1D Tensor", "Sparse indices vector.") + .add_argument("sparse_indptr", "1D Tensor", "Sparse index pointer vector.") .set_support_level(1) .add_type_rel("SparseAdd", SparseAddRel); From 165f15bef4692d45e63da8ff5a476a88f9f998cc Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Thu, 18 Feb 2021 23:10:08 +0530 Subject: [PATCH 6/9] [2] Review comment handled --- python/tvm/relay/op/nn/nn.py | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/python/tvm/relay/op/nn/nn.py b/python/tvm/relay/op/nn/nn.py index 2b854c840dc4..f60d72cc02ff 100644 --- a/python/tvm/relay/op/nn/nn.py +++ b/python/tvm/relay/op/nn/nn.py @@ -2174,6 +2174,21 @@ def sparse_add(dense_mat, sparse_mat): ------- result: tvm.relay.Expr The computed result. + + Examples + ------- + .. code-block:: python + dense_data = [[ 3., 4., 4. ] + [ 4., 2., 5. ]] + sparse_data = [4., 8.] + sparse_indices =[0, 2] + sparse_indptr =[0, 1, 2] + dense_shape = [2, 3] + + output = relay.sparse_add(dense_data, sparse_data, sparse_indices, sparse_indptr) + + output = [[ 7., 4., 4. ] + [ 4., 2., 13. ]] """ if hasattr(sparse_mat, "indices"): return _make.sparse_add(dense_mat, sparse_mat.data, sparse_mat.indices, sparse_mat.indptr) From 5286226886879ce45b54d0f6cce6fb2922f0c267 Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Wed, 24 Feb 2021 21:50:17 +0530 Subject: [PATCH 7/9] [3] Review comment handled --- python/tvm/relay/op/nn/nn.py | 5 ++--- src/relay/op/nn/sparse.cc | 8 +++++--- 2 files changed, 7 insertions(+), 6 deletions(-) diff --git a/python/tvm/relay/op/nn/nn.py b/python/tvm/relay/op/nn/nn.py index f60d72cc02ff..2f6c898f4c70 100644 --- a/python/tvm/relay/op/nn/nn.py +++ b/python/tvm/relay/op/nn/nn.py @@ -2179,16 +2179,15 @@ def sparse_add(dense_mat, sparse_mat): ------- .. code-block:: python dense_data = [[ 3., 4., 4. ] - [ 4., 2., 5. ]] + [ 4., 2., 5. ]] sparse_data = [4., 8.] sparse_indices =[0, 2] sparse_indptr =[0, 1, 2] - dense_shape = [2, 3] output = relay.sparse_add(dense_data, sparse_data, sparse_indices, sparse_indptr) output = [[ 7., 4., 4. ] - [ 4., 2., 13. ]] + [ 4., 2., 13. ]] """ if hasattr(sparse_mat, "indices"): return _make.sparse_add(dense_mat, sparse_mat.data, sparse_mat.indices, sparse_mat.indptr) diff --git a/src/relay/op/nn/sparse.cc b/src/relay/op/nn/sparse.cc index 94d088d08e61..406c00683eff 100644 --- a/src/relay/op/nn/sparse.cc +++ b/src/relay/op/nn/sparse.cc @@ -202,10 +202,12 @@ bool SparseAddRel(const Array& types, int num_inputs, const Attrs& attrs, ICHECK_EQ(types.size(), 5); const auto* dense_data = types[0].as(); const auto* sparse_data = types[1].as(); - ICHECK(reporter->Assert(sparse_data->dtype == dense_data->dtype)); - ICHECK(reporter->Assert(sparse_data->shape.size() == 1)); + ICHECK(reporter->Assert(sparse_data->dtype == dense_data->dtype)) + << "sparse tensor and dense tensor datatype should match."; + ICHECK(reporter->Assert(sparse_data->shape.size() == 1)) << "sparse data tensor should be 1D."; const auto* sparse_indices = types[2].as(); - ICHECK(reporter->Assert(sparse_indices->shape.size() == 1)); + ICHECK(reporter->Assert(sparse_indices->shape.size() == 1)) + << "sparse indices tensor should be 1D."; reporter->Assign(types[4], TensorType(dense_data->shape, dense_data->dtype)); return true; From d183828dc1ee38ddf7016bba2b7dd4261783c5a4 Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Thu, 4 Mar 2021 21:11:34 +0530 Subject: [PATCH 8/9] [4] Review comment handled --- python/tvm/relay/op/nn/nn.py | 4 +- python/tvm/relay/op/strategy/generic.py | 7 --- python/tvm/topi/generic/nn.py | 17 ------- python/tvm/topi/nn/sparse.py | 16 +++---- src/relay/op/nn/sparse.cc | 2 +- .../frontend/tensorflow/test_forward.py | 14 +++--- tests/python/topi/python/test_topi_sparse.py | 46 ++++++++++--------- 7 files changed, 43 insertions(+), 63 deletions(-) diff --git a/python/tvm/relay/op/nn/nn.py b/python/tvm/relay/op/nn/nn.py index 2f6c898f4c70..62882e401344 100644 --- a/python/tvm/relay/op/nn/nn.py +++ b/python/tvm/relay/op/nn/nn.py @@ -2165,10 +2165,10 @@ def sparse_add(dense_mat, sparse_mat): Parameters ---------- dense_mat : tvm.relay.Expr - The input dense matrix for the matrix multiplication + The input dense matrix for the matrix addition sparse_mat : Union[namedtuple, Tuple[ndarray, ndarray, ndarray]]. - The input sparse matrix for the matrix multiplication. + The input sparse matrix(CSR) for the matrix addition. Returns ------- diff --git a/python/tvm/relay/op/strategy/generic.py b/python/tvm/relay/op/strategy/generic.py index f2775d07ea7c..04f25640574a 100644 --- a/python/tvm/relay/op/strategy/generic.py +++ b/python/tvm/relay/op/strategy/generic.py @@ -822,13 +822,6 @@ def sparse_add_strategy(attrs, inputs, out_type, target): return strategy -@generic_func -def schedule_sparse_add(attrs, outs, target): - """schedule sparse_add""" - with target: - return topi.generic.schedule_sparse_add(outs) - - # sparse_transpose @generic_func def schedule_sparse_transpose(attrs, outs, target): diff --git a/python/tvm/topi/generic/nn.py b/python/tvm/topi/generic/nn.py index 49281e0356a6..60ccd0d36abf 100644 --- a/python/tvm/topi/generic/nn.py +++ b/python/tvm/topi/generic/nn.py @@ -730,23 +730,6 @@ def schedule_sparse_transpose(outs): return _default_schedule(outs, False) -def schedule_sparse_add(outs): - """Schedule for sparse_add - - Parameters - ---------- - outs: Array of Tensor - The computation graph description of sparse_add - in the format of an array of tensors. - - Returns - ------- - sch: Schedule - The computation schedule for the op. - """ - return _default_schedule(outs, False) - - def schedule_batch_matmul(outs): """Schedule for batch_matmul diff --git a/python/tvm/topi/nn/sparse.py b/python/tvm/topi/nn/sparse.py index de46117fa811..756110624aa1 100644 --- a/python/tvm/topi/nn/sparse.py +++ b/python/tvm/topi/nn/sparse.py @@ -477,16 +477,16 @@ def sparse_add(dense_data, sparse_data, sparse_indices, sparse_indptr): Parameters ---------- dense_data : tvm.te.Tensor - 2-D with shape [M, N], float32 + 2-D with shape [M, N] sparse_data : tvm.te.Tensor - 1-D with shape [nnz] (CSR) or + 1-D with shape [nnz] (CSR) sparse_indices : tvm.te.Tensor - 1-D with shape [nnz] (CSR) or + 1-D with shape [nnz] (CSR) sparse_indptr : tvm.te.Tensor - 1-D with shape [M + 1] (CSR) or + 1-D with shape [M + 1] (CSR) Returns ------- @@ -510,11 +510,11 @@ def _csr_add_ir(dense_data, sparse_data, sparse_indices, sparse_indptr, out_data out_data_ptr = irb.buffer_ptr(out_data) - with irb.for_range(0, oshape[0], kind="serial", name="row") as row: - with irb.for_range(0, oshape[1], kind="serial", name="col") as col: + with irb.for_range(0, oshape[0], kind="vectorize", name="row") as row: + with irb.for_range(0, oshape[1], kind="parallel", name="col") as col: out_data_ptr[row, col] = dense_data_ptr[row, col] - with irb.for_range(0, oshape[0], kind="serial", name="row") as row: + with irb.for_range(0, oshape[0], kind="parallel", name="row") as row: offset = sparse_indptr_ptr[row] diff = sparse_indptr_ptr[row + 1] - sparse_indptr_ptr[row] with irb.for_range(0, diff, kind="serial", name="idx") as idx: @@ -535,5 +535,5 @@ def _csr_add_ir(dense_data, sparse_data, sparse_indices, sparse_indptr, out_data sparse_indices_inp.dtype, sparse_indptr_inp.dtype, ], - name="out", + name="sparse_add_csr_output", ) diff --git a/src/relay/op/nn/sparse.cc b/src/relay/op/nn/sparse.cc index 406c00683eff..b1a16f18b623 100644 --- a/src/relay/op/nn/sparse.cc +++ b/src/relay/op/nn/sparse.cc @@ -199,7 +199,7 @@ RELAY_REGISTER_OP("nn.sparse_transpose") // relay.nn.sparse_add bool SparseAddRel(const Array& types, int num_inputs, const Attrs& attrs, const TypeReporter& reporter) { - ICHECK_EQ(types.size(), 5); + ICHECK_EQ(types.size(), 5) << "expecting 4 inputs and 1 output."; const auto* dense_data = types[0].as(); const auto* sparse_data = types[1].as(); ICHECK(reporter->Assert(sparse_data->dtype == dense_data->dtype)) diff --git a/tests/python/frontend/tensorflow/test_forward.py b/tests/python/frontend/tensorflow/test_forward.py index 6b0b13a6bc73..fa27dee37699 100644 --- a/tests/python/frontend/tensorflow/test_forward.py +++ b/tests/python/frontend/tensorflow/test_forward.py @@ -2364,7 +2364,9 @@ def _test_sparse_add(indices, values, A_shape, B_shape, dtype, flip=False): # TODO(ANSHUMAN87): support both sparse input case with tf.Graph().as_default(): - A_sp = tf.sparse.SparseTensor(indices=indices, values=values, dense_shape=A_shape) + A_sp = tf.sparse.SparseTensor( + indices=indices, values=np.array(values).astype(dtype), dense_shape=A_shape + ) B = tf.placeholder(shape=B_shape, dtype=dtype, name="B") # TODO(ANSHUMAN87): support user input threashold values @@ -2391,11 +2393,11 @@ def test_sparse_add(): # [0, 0, 0, 0]] # # ------------------------------------------------------------------ - - _test_sparse_add([[0, 0], [1, 2]], [4.0, 8.0], [3, 4], [3, 4], "float32") - _test_sparse_add([[0, 0], [1, 2]], [4.0, 8.0], [3, 4], [3, 4], "float32", True) - _test_sparse_add([[0, 0], [1, 3], [4, 3]], [3.0, 6.0, 9.0], [5, 5], [5, 5], "float32") - _test_sparse_add([[0, 0], [1, 3], [4, 3]], [3.0, 6.0, 9.0], [5, 5], [5, 5], "float32", True) + for dtype_inp in ["float32", "float64", "int32"]: + _test_sparse_add([[0, 0], [1, 2]], [4.0, 8.0], [3, 4], [3, 4], dtype_inp) + _test_sparse_add([[0, 0], [1, 2]], [4.0, 8.0], [3, 4], [3, 4], dtype_inp, True) + _test_sparse_add([[0, 0], [1, 3], [4, 3]], [3.0, 6.0, 9.0], [5, 5], [5, 5], dtype_inp) + _test_sparse_add([[0, 0], [1, 3], [4, 3]], [3.0, 6.0, 9.0], [5, 5], [5, 5], dtype_inp, True) ####################################################################### diff --git a/tests/python/topi/python/test_topi_sparse.py b/tests/python/topi/python/test_topi_sparse.py index 5d92694cfaa4..d84bd1530587 100644 --- a/tests/python/topi/python/test_topi_sparse.py +++ b/tests/python/topi/python/test_topi_sparse.py @@ -527,28 +527,30 @@ def test_sparse_dense_padded_alter_op(): def test_sparse_add_csr(): - M, K, density = 3, 49, 0.2 - X_np = np.random.randn(M, K).astype("float32") - Y_sp_np = sp.random(M, K, density=density, format="csr", dtype="float32") - Y_np = Y_sp_np.todense() - Z_np = X_np + Y_np - - Y_data = te.placeholder(shape=Y_sp_np.data.shape, dtype=str(Y_sp_np.data.dtype)) - Y_indices = te.placeholder(shape=Y_sp_np.indices.shape, dtype=str(Y_sp_np.indices.dtype)) - Y_indptr = te.placeholder(shape=Y_sp_np.indptr.shape, dtype=str(Y_sp_np.indptr.dtype)) - X = te.placeholder(shape=X_np.shape, dtype=str(X_np.dtype)) - Z = topi.nn.sparse_add(X, Y_data, Y_indices, Y_indptr) - s = te.create_schedule(Z.op) - func = tvm.build(s, [X, Y_data, Y_indices, Y_indptr, Z]) - Z_tvm = tvm.nd.array(np.zeros(Z_np.shape, dtype=Z_np.dtype)) - func( - tvm.nd.array(X_np), - tvm.nd.array(Y_sp_np.data), - tvm.nd.array(Y_sp_np.indices), - tvm.nd.array(Y_sp_np.indptr), - Z_tvm, - ) - tvm.testing.assert_allclose(Z_tvm.asnumpy(), Z_np, atol=1e-4, rtol=1e-4) + for indices_dtype in ["int32", "int64"]: + for data_dtype in ["float32", "float64"]: + M, K, density = 3, 49, 0.2 + X_np = np.random.randn(M, K).astype(data_dtype) + Y_sp_np = sp.random(M, K, density=density, format="csr", dtype=data_dtype) + Y_np = Y_sp_np.todense() + Z_np = X_np + Y_np + + Y_data = te.placeholder(shape=Y_sp_np.data.shape, dtype=data_dtype) + Y_indices = te.placeholder(shape=Y_sp_np.indices.shape, dtype=indices_dtype) + Y_indptr = te.placeholder(shape=Y_sp_np.indptr.shape, dtype=indices_dtype) + X = te.placeholder(shape=X_np.shape, dtype=data_dtype) + Z = topi.nn.sparse_add(X, Y_data, Y_indices, Y_indptr) + s = te.create_schedule(Z.op) + func = tvm.build(s, [X, Y_data, Y_indices, Y_indptr, Z]) + Z_tvm = tvm.nd.array(np.zeros(Z_np.shape, dtype=Z_np.dtype)) + func( + tvm.nd.array(X_np.astype(data_dtype)), + tvm.nd.array(Y_sp_np.data.astype(data_dtype)), + tvm.nd.array(Y_sp_np.indices.astype(indices_dtype)), + tvm.nd.array(Y_sp_np.indptr.astype(indices_dtype)), + Z_tvm, + ) + tvm.testing.assert_allclose(Z_tvm.asnumpy(), Z_np, atol=1e-4, rtol=1e-4) if __name__ == "__main__": From 0e0b19c88b6e3ccae67e28a2305cd5422d5ad281 Mon Sep 17 00:00:00 2001 From: ANSHUMAN TRIPATHY Date: Tue, 9 Mar 2021 12:05:15 +0530 Subject: [PATCH 9/9] [5] Review comment handled --- python/tvm/relay/op/nn/nn.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/tvm/relay/op/nn/nn.py b/python/tvm/relay/op/nn/nn.py index 62882e401344..a1147fec4d7e 100644 --- a/python/tvm/relay/op/nn/nn.py +++ b/python/tvm/relay/op/nn/nn.py @@ -2152,7 +2152,7 @@ def sparse_transpose(x): def sparse_add(dense_mat, sparse_mat): r""" Computes the matrix addition of `dense_mat` and `sparse_mat`, where `dense_mat` is - a dense matrix and `sparse_mat` is a sparse (either BSR or CSR) namedtuple with + a dense matrix and `sparse_mat` is a sparse (CSR) namedtuple with fields `data`, `indices`, and `indptr`. .. math::