masahi commented on a change in pull request #7334:
URL: https://github.com/apache/tvm/pull/7334#discussion_r564021971
##########
File path: python/tvm/topi/cuda/scan.py
##########
@@ -251,99 +263,98 @@ def scan_thrust(data, output_dtype, exclusive=True,
return_reduction=False):
Whether or not do exclusive or inclusive scan.
return_reduction: bool, optional
- Whether or not return a 1-D tensor storing the reduction of each row.
+ Whether or not return a (N-1)-D tensor storing the reduction of each
scan axis.
Reductions are computed as part of the upsweep pass, so there is no
extra cost.
- If False, reductions are ignored.
+ If False, reductions are ignored. It must be False when exclusive is
False.
+
+ biop: string, optional
+ A string specifying which binary operator to use. Currently only "sum"
is supported.
Returns
-------
output : tvm.te.Tensor
- 1-D tensor that is the exclusive scan of the input, or
- 2-D tensor storing the exclusive scan of each row.
+ A N-D tensor of the same rank N and shape as the input data.
reduction : tvm.te.Tensor, optional
- 1-D tensor storing the reduction of each row.
+ (N-1)-D tensor storing the reduction of each scan axis.
Returned if return_reduction is True.
"""
data_buf = tvm.tir.decl_buffer(data.shape, data.dtype, "data_buf",
data_alignment=8)
output_buf = tvm.tir.decl_buffer(data.shape, output_dtype, "output_buf",
data_alignment=8)
+ binop_to_thrust_func_name = {"sum": "tvm.contrib.thrust.sum_scan"}
output = te.extern(
[data.shape],
[data],
lambda ins, outs: tvm.tir.call_packed(
- "tvm.contrib.thrust.sum_scan", ins[0], outs[0], exclusive
+ binop_to_thrust_func_name[binop], ins[0], outs[0], exclusive
),
dtype=[output_dtype],
in_buffers=[data_buf],
out_buffers=[output_buf],
- name="exclusive_sum_scan2d",
- tag="exclusive_sum_scan2d_gpu",
+ name="exclusive_scan_thrust",
+ tag="exclusive_scan_thrust_gpu",
)
if return_reduction:
assert exclusive, "return_reduction should be False for inclusive scan"
- reduction = get_reduction_from_exclusive_scan(data, output)
+ reduction = get_reduction_from_exclusive_scan(data, output, binop)
return output, reduction
return output
-def exclusive_scan(data, axis=-1, return_reduction=False, output_dtype=None):
- """Do exclusive scan on 1D input or along rows of 2D input.
+def exclusive_scan(data, axis=-1, return_reduction=False, output_dtype=None,
binop="sum"):
+ """Do exclusive scan on 1D or multidimensional input.
Parameters
----------
data : tvm.te.Tensor
- Input data. 1-D tensor with shape [scan_axis_size], or
- 2-D tensor with shape [batch_size, scan_axis_size].
+ Input data of any shape.
axis: int, optional
- The axis to do scan on. For now, only the inner most axis is supported.
+ The axis to do scan on. By default, scan is done on the innermost axis.
return_reduction: bool, optional
- Whether or not return a 1-D tensor storing the reduction of each row.
+ Whether or not return a tensor storing the reduction over each scan
axis.
+ If the input rank is N, this tensor is of rank N - 1.
Reductions are computed as part of the upsweep pass, so there is no
extra cost.
If False, reductions are ignored.
output_dtype: string, optional
The dtype of the output scan tensor. If not provided, the dtype of the
input is used.
+ biop: string, optional
+ A string specifying which binary operator to use. Currently only "sum"
is supported.
+
Returns
-------
output : tvm.te.Tensor
- 1-D tensor that is the exclusive scan of the input, or
- 2-D tensor storing the exclusive scan of each row.
+ A N-D tensor of the same rank N and shape as the input data.
reduction : tvm.te.Tensor, optional
- 1-D tensor storing the reduction of each row.
+ (N-1)-D tensor storing the reduction of each scan axis.
Returned if return_reduction is True.
"""
- # TODO(masahi): Support other binary operators
- ndim = len(data.shape)
- if axis < 0:
- axis += ndim
- assert axis == ndim - 1, "Only support scan on the inner most axis."
-
- if output_dtype is None:
- output_dtype = data.dtype
- target = tvm.target.Target.current()
- if target and target.kind.name == "cuda" and is_thrust_available():
- return scan_thrust(data, output_dtype, exclusive=True,
return_reduction=return_reduction)
+ def do_scan(data, output_dtype):
+ target = tvm.target.Target.current()
+ if target and target.kind.name == "cuda" and is_thrust_available():
Review comment:
It's a bit tricky, since `exclusive_scan` is called by other ops, I need
to introduce separate implementation and strategy for every op that uses it.
Currently they are `get_valid_counts`, `argwhere` and `cumsum`. Soon I'll add
`unique`.
----------------------------------------------------------------
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]
