def upsampling(data,
scale_h,
scale_w,
layout="NCHW",
method='nearest_neighbor',
align_corners=False):
"""Perform upsampling on the data.
Nearest neighbor and bilinear upsampling are supported.
Parameters
----------
inputs : tvm.te.Tensor
inputs is a 4-D tensor with shape
[batch, channel, in_height, in_width]
or [batch, in_height, in_width, channel]
scale_h : float
Scaling factor for height
scale_w : float
Scaling factor for width
layout : string, optional
either "NCHW" or "NHWC"
method : {"bilinear", "nearest_neighbor", "bicubic"}
Method to be used for upsampling.
Returns
-------
output : tvm.te.Tensor
4-D with shape [batch, channel, in_height*scale_h, in_width*scale_w]
or [batch, in_height*scale, in_width*scale, channel]
"""
base_layout = layout[0:4]
if base_layout == "NCHW":
out_shape = (simplify(
topi.cast(te.round(data.shape[2] * scale_h), data.shape[2].dtype)),
simplify(
topi.cast(te.round(data.shape[3] * scale_w),
data.shape[3].dtype)))
elif layout == "NHWC":
out_shape = (simplify(
topi.cast(te.round(data.shape[1] * scale_h), data.shape[1].dtype)),
simplify(
topi.cast(te.round(data.shape[2] * scale_w),
data.shape[2].dtype)))
else:
raise ValueError("not support this layout {} yet".format(layout))
coord_trans = "align_corners" if align_corners else "asymmetric"
return topi.image.resize(data,
out_shape,
layout=layout,
method=method,
coordinate_transformation_mode=coord_trans)
def _pool(i, c, ph, pw):
roi = rois[i]
batch_index = roi[0].astype("int32")
roi_start_w, roi_start_h, roi_end_w, roi_end_h = roi[1], roi[2], roi[
3], roi[4]
roi_start_h = te.round(roi_start_h * spatial_scale).astype("int32")
roi_start_w = te.round(roi_start_w * spatial_scale).astype("int32")
roi_end_h = te.round(roi_end_h * spatial_scale).astype("int32")
roi_end_w = te.round(roi_end_w * spatial_scale).astype("int32")
# force malformed ROIs to be 1x1
roi_h = tvm.te.max(roi_end_h - roi_start_h + 1,
tvm.tir.const(1, "int32"))
roi_w = tvm.te.max(roi_end_w - roi_start_w + 1,
tvm.tir.const(1, "int32"))
bin_h = roi_h.astype(dtype) / pooled_size_h
bin_w = roi_w.astype(dtype) / pooled_size_w
# use epsilon to prevent floating point precision loss in floor/ceil
epsilon = tvm.tir.const(0.00001, dtype)
hstart = te.floor(ph * bin_h + epsilon).astype("int32")
wstart = te.floor(pw * bin_w + epsilon).astype("int32")
hend = te.ceil((ph + 1) * bin_h - epsilon).astype("int32")
wend = te.ceil((pw + 1) * bin_w - epsilon).astype("int32")
hstart = tvm.te.min(tvm.te.max(hstart + roi_start_h, 0), height)
wstart = tvm.te.min(tvm.te.max(wstart + roi_start_w, 0), width)
hend = tvm.te.min(tvm.te.max(hend + roi_start_h, 0), height)
wend = tvm.te.min(tvm.te.max(wend + roi_start_w, 0), width)
non_empty = tvm.tir.all(hstart < hend, wstart < wend)
min_value = lambda dtype: tvm.tir.if_then_else(
non_empty, tvm.te.min_value(dtype), tvm.tir.const(0.0, dtype))
# pylint: disable=unnecessary-lambda
_max = te.comm_reducer(lambda x, y: tvm.te.max(x, y),
min_value,
name="max")
rh = te.reduce_axis((0, hend - hstart), "rh")
rw = te.reduce_axis((0, wend - wstart), "rw")
return _max(data[batch_index, c, hstart + rh, wstart + rw],
axis=[rh, rw])
def _nearest_neighbor(*indices):
n, c, z, y, x, cc = _get_indices(*indices)
in_z = z_ratio * z
in_y = y_ratio * y
in_x = x_ratio * x
if coordinate_transformation_mode == "align_corners":
zint = te.round(in_z).astype('int32')
yint = te.round(in_y).astype('int32')
xint = te.round(in_x).astype('int32')
elif coordinate_transformation_mode in ["asymmetric", "half_pixel"]:
# Add epsilon to floor to prevent gpu rounding errors.
epsilon = 1e-5
zint = te.floor(in_z + epsilon).astype('int32')
yint = te.floor(in_y + epsilon).astype('int32')
xint = te.floor(in_x + epsilon).astype('int32')
else:
raise ValueError("Unsupported coordinate_transformation_mode: {}".format(
coordinate_transformation_mode))
return _cast_output(_get_pixel(n, c, zint, yint, xint, cc))
def round(x):
"""Round elements of x to nearest integer.
Parameters
----------
x : tvm.te.Tensor
Input argument.
Returns
-------
y : tvm.te.Tensor
The result.
"""
return te.compute(x.shape, lambda *i: te.round(x(*i)))
def test_const_fold4():
x1 = tvm.tir.const(4, "int32")
x2 = x1 + 5
tdiv = tvm.tir.truncdiv
assert isinstance(x2, tvm.tir.IntImm) and x2.value == 9
x3 = tdiv(x2, 3)
assert isinstance(x3, tvm.tir.IntImm) and x3.value == 3
x4 = x3 + 0.55
assert isinstance(x4, tvm.tir.FloatImm) and abs(x4.value - 3.55) < 1e-6
x5 = te.ceil(x4)
assert isinstance(x5, tvm.tir.FloatImm) and x5.value == 4
x6 = x5.astype("int")
assert isinstance(x6, tvm.tir.IntImm) and x6.value == 4, "x6={}".format(x6)
y = (te.round((tvm.tir.const(6.5, "float32") - 1) / 1.5) + 2).astype("int")
assert isinstance(y, tvm.tir.IntImm) and y.value == 6
def _compute_intn(dtype, value, *indices):
assert output_scale is not None and output_zero_point is not None
const_min = tvm.tir.min_value(dtype)
const_max = tvm.tir.max_value(dtype)
# Use indexmod to handle both scalar and per-channel QNN parameters.
scale_idx = tir.indexmod(indices[axis], topi.shape(output_scale)[0])
zp_idx = tir.indexmod(indices[axis], topi.shape(output_zero_point)[0])
return te.max(
te.min(
te.round(value[indices] / output_scale[scale_idx]) +
output_zero_point[zp_idx],
const_max,
),
const_min,
)
def upsampling3d(data, scale_d, scale_h, scale_w, layout="NCDHW", method='nearest_neighbor',
coordinate_transformation_mode="half_pixel"):
"""Perform upsampling on the data.
Nearest neighbor and bilinear upsampling are supported.
Parameters
----------
inputs : tvm.te.Tensor
inputs is a 5-D tensor with shape
[batch, channel, in_depth, in_height, in_width]
or [batch, in_depth, in_height, in_width, channel]
scale_d : float
Scaling factor for depth
scale_h : float
Scaling factor for height
scale_w : float
Scaling factor for width
layout : string, optional
either "NCDHW" or "NDHWC"
method : {"trilinear", "nearest_neighbor"}
Method to be used for upsampling.
coordinate_transformation_mode: string, optional
Describes how to transform the coordinate in the resized tensor
to the coordinate in the original tensor.
Refer to the ONNX Resize operator specification for details.
Available options are "half_pixel", "align_corners" and "asymmetric".
Returns
-------
output : tvm.te.Tensor
5-D with shape [batch, channel, in_depth*scale, in_height*scale, in_width*scale]
or [batch, in_depth*scale, in_height*scale, in_width*scale, channel]
"""
base_layout = layout[0:5]
if base_layout == "NCDHW":
out_shape = (simplify(topi.cast(te.round(data.shape[2] * scale_d), data.shape[2].dtype)),
simplify(topi.cast(te.round(data.shape[3] * scale_h), data.shape[3].dtype)),
simplify(topi.cast(te.round(data.shape[4] * scale_w), data.shape[4].dtype)))
elif layout == "NDHWC":
out_shape = (simplify(topi.cast(te.round(data.shape[1] * scale_d), data.shape[1].dtype)),
simplify(topi.cast(te.round(data.shape[2] * scale_h), data.shape[2].dtype)),
simplify(topi.cast(te.round(data.shape[3] * scale_w), data.shape[3].dtype)))
else:
raise ValueError("not support this layout {} yet".format(layout))
return topi.image.resize3d(data, out_shape, layout=layout, method=method,
coordinate_transformation_mode=coordinate_transformation_mode)
def get_closest_index(in_x, rounding_method, boxes):
"""get the closest index to a value based on a certain rounding method"""
if rounding_method == "round" or boxes is not None:
closest_x_index = te.round(in_x).astype("int32")
elif rounding_method == "round_prefer_floor":
closest_x_index = te.ceil(in_x - 0.5).astype("int32")
elif rounding_method == "round_prefer_ceil":
closest_x_index = te.floor(in_x + 0.5).astype("int32")
elif rounding_method == "floor":
# Add epsilon to floor to prevent gpu rounding errors.
epsilon = 1e-5
closest_x_index = te.floor(in_x + epsilon).astype("int32")
elif rounding_method == "ceil":
# Subract epsilon from ceil to prevent gpu rounding errors.
epsilon = 1e-5
closest_x_index = te.ceil(in_x - epsilon).astype("int32")
else:
raise ValueError("Uknown rounding method: {}".format(rounding_method))
return closest_x_index
def upsampling(data,
scale_h,
scale_w,
layout="NCHW",
method='nearest_neighbor',
align_corners=False,
output_shape=None):
"""Perform upsampling on the data.
Nearest neighbor and bilinear upsampling are supported.
Parameters
----------
inputs : tvm.te.Tensor
inputs is a 4-D tensor with shape
[batch, channel, in_height, in_width]
or [batch, in_height, in_width, channel]
scale_h : float
Scaling factor for height
scale_w : float
Scaling factor for width
layout : string, optional
either "NCHW" or "NHWC"
method : {"bilinear", "nearest_neighbor", "bicubic"}
Method to be used for upsampling.
output_shape: tvm.tir.container.Array, optional
Shape to return. If left None will be inferred
(If shape is determined dynamically, pass out_dtype.shape as output_shape)
Returns
-------
output : tvm.te.Tensor
4-D with shape [batch, channel, in_height*scale_h, in_width*scale_w]
or [batch, in_height*scale, in_width*scale, channel]
"""
base_layout = layout[0:4]
if base_layout == "NCHW":
if not output_shape: #static case
scaled_h = data.shape[2] * scale_h
scaled_w = data.shape[3] * scale_w
reshape_size = (simplify(
topi.cast(te.round(scaled_h), data.shape[2].dtype)),
simplify(
topi.cast(te.round(scaled_w),
data.shape[3].dtype)))
else: #dynamic case -- we don't need to scale; already done in shape func
reshape_size = (simplify(
topi.cast(te.round(output_shape[2]), output_shape[2].dtype)),
simplify(
topi.cast(te.round(output_shape[3]),
output_shape[3].dtype)))
elif layout == "NHWC":
if not output_shape: #static case
scaled_h = data.shape[1] * scale_h
scaled_w = data.shape[2] * scale_w
reshape_size = (simplify(
topi.cast(te.round(scaled_h), data.shape[1].dtype)),
simplify(
topi.cast(te.round(scaled_w),
data.shape[2].dtype)))
else: #dynamic case
reshape_size = (simplify(
topi.cast(te.round(output_shape[1]), output_shape[1].dtype)),
simplify(
topi.cast(te.round(output_shape[2]),
output_shape[2].dtype)))
else:
raise ValueError("not support this layout {} yet".format(layout))
coord_trans = "align_corners" if align_corners else "asymmetric"
return topi.image.resize(data,
reshape_size,
layout=layout,
method=method,
coordinate_transformation_mode=coord_trans,
output_shape=output_shape)
def upsampling3d(data,
scale_d,
scale_h,
scale_w,
layout="NCDHW",
method='nearest_neighbor',
coordinate_transformation_mode="half_pixel",
output_shape=None):
"""Perform upsampling on the data.
Nearest neighbor and bilinear upsampling are supported.
Parameters
----------
inputs : tvm.te.Tensor
inputs is a 5-D tensor with shape
[batch, channel, in_depth, in_height, in_width]
or [batch, in_depth, in_height, in_width, channel]
scale_d : float
Scaling factor for depth
scale_h : float
Scaling factor for height
scale_w : float
Scaling factor for width
layout : string, optional
either "NCDHW" or "NDHWC"
method : {"trilinear", "nearest_neighbor"}
Method to be used for upsampling.
coordinate_transformation_mode: string, optional
Describes how to transform the coordinate in the resized tensor
to the coordinate in the original tensor.
Refer to the ONNX Resize operator specification for details.
Available options are "half_pixel", "align_corners" and "asymmetric".
output_shape: tvm.tir.container.Array, optional
Shape to return. If left None will be inferred
(If shape is determined dynamically, pass out_dtype.shape as output_shape)
Returns
-------
output : tvm.te.Tensor
5-D with shape [batch, channel, in_depth*scale, in_height*scale, in_width*scale]
or [batch, in_depth*scale, in_height*scale, in_width*scale, channel]
"""
base_layout = layout[0:5]
if base_layout == "NCDHW":
if not output_shape: # static case
scaled_d = data.shape[2] * scale_d
scaled_h = data.shape[3] * scale_h
scaled_w = data.shape[4] * scale_w
resize_shape = (simplify(
topi.cast(te.round(scaled_d), data.shape[2].dtype)),
simplify(
topi.cast(te.round(scaled_h),
data.shape[3].dtype)),
simplify(
topi.cast(te.round(scaled_w),
data.shape[4].dtype)))
else: # dynamic case -- don't need to scale; already done in shape func
resize_shape = (simplify(
topi.cast(te.round(output_shape[2]), data.shape[2].dtype)),
simplify(
topi.cast(te.round(output_shape[3]),
data.shape[3].dtype)),
simplify(
topi.cast(te.round(output_shape[4]),
data.shape[4].dtype)))
elif layout == "NDHWC":
if not output_shape: # static case
scaled_d = data.shape[1] * scale_d
scaled_h = data.shape[2] * scale_h
scaled_w = data.shape[3] * scale_w
resize_shape = (simplify(
topi.cast(te.round(scaled_d), data.shape[1].dtype)),
simplify(
topi.cast(te.round(scaled_h),
data.shape[2].dtype)),
simplify(
topi.cast(te.round(scaled_w),
data.shape[3].dtype)))
else: # dynamic case
resize_shape = (simplify(
topi.cast(te.round(output_shape[1]), data.shape[1].dtype)),
simplify(
topi.cast(te.round(output_shape[2]),
data.shape[2].dtype)),
simplify(
topi.cast(te.round(output_shape[3]),
data.shape[3].dtype)))
else:
raise ValueError("not support this layout {} yet".format(layout))
return topi.image.resize3d(
data,
resize_shape,
layout=layout,
method=method,
coordinate_transformation_mode=coordinate_transformation_mode)
def test_basic_operation():
np.random.seed(0)
shape = (10, 10)
x = te.var("x", dtype='float32')
k = te.reduce_axis((0, 10), name="k")
l = te.reduce_axis((0, 10), name="l")
A0 = te.placeholder(shape, name='A0')
A1 = te.placeholder(shape, name='A1')
zeros = np.zeros(shape)
B = te.compute(shape, lambda i, j: A0[i, j], name='B')
check_grad(B, [A0])
B = te.compute(shape, lambda i, j: A0[i, j] + A1[i, j], name='B')
check_grad(B, [A0, A1])
B = te.compute(shape, lambda i, j: A0[i, j] + A0[j, i], name='B')
check_grad(B, A0)
B = te.compute(shape, lambda i, j: te.floor(A0[i, j]), name='B')
check_grad(B, A0, desired_grads=[zeros])
B = te.compute(shape, lambda i, j: te.ceil(A0[i, j]), name='B')
check_grad(B, A0, desired_grads=[zeros])
B = te.compute(shape, lambda i, j: te.trunc(A0[i, j]), name='B')
check_grad(B, A0, desired_grads=[zeros])
B = te.compute(shape, lambda i, j: te.round(A0[i, j]), name='B')
check_grad(B, A0, desired_grads=[zeros])
B = te.compute(shape, lambda i, j: A0[i, j] + te.exp(A0[j, i]), name='B')
check_grad(B, A0)
B = te.compute(
shape,
lambda i, j: te.log(0.1 + te.abs(A0[i, j] + te.exp(A0[j, i]))),
name='B')
check_grad(B, A0)
B = te.compute(shape,
lambda i, j: te.sigmoid(A0[i, j] * A0[i, j] * A0[j, i]),
name='B')
check_grad(B, A0)
B = te.compute(shape,
lambda i, j: te.tanh(A0[i, j] * A0[i, j] * A0[j, i]),
name='B')
check_grad(B, A0)
B = te.compute(shape,
lambda i, j: te.sqrt(A0[i, j] * A0[i, j] * A0[j, i]),
name='B')
check_grad(B, A0, data_range=(0.1, 10))
B = te.compute(shape,
lambda i, j: te.power(te.abs(A0[i, j]), A0[j, i]),
name='B')
check_grad(B, A0, data_range=(-4, 4))
B = te.compute(shape, lambda i, j: A0[i, j] * A0[j, i], name='B')
check_grad(B, A0)
B = te.compute((10, ),
lambda i: te.sum(A0[i, k] * A0[k, i], axis=k),
name='B')
check_grad(B, A0)
B = te.compute(shape,
lambda i, j: te.sum(A0[i, k] * A0[k, i] + 5, axis=k),
name='B')
check_grad(B, A0)
B = te.compute(shape,
lambda i, j: te.max(A0[i, k] * A0[k, j] + 5, axis=k),
name='B')
check_grad(B, A0)
B = te.compute(shape,
lambda i, j: A0[i, j] * (A1[j, i] + A0[j, i]),
name='B')
check_grad(B, [A0, A1])
B = te.compute(shape,
lambda i, j: te.sum(
A0[k, k] - A0[te.min(j + k, 9), j] * A0[i, k], axis=k),
name='B')
check_grad(B, A0)
def fcombine(x, y):
return x * y
def fidentity(t0):
return tvm.tir.const(1, t0)
prod = te.comm_reducer(fcombine, fidentity, name='prod')
B = te.compute((10, 10),
lambda i, j: prod(A0[i, k] + A0[k, i], axis=k),
name='B')
check_grad(B, A0)
X = te.placeholder((10, ), name='X')
A = te.compute((10, ), lambda i: X[i] + X[9 - i])
B = te.compute((10, ), lambda i: X[i] * X[9 - i])
Y = topi.tensordot(A, B, 1)
check_grad(Y, X)
def resize_nearest_neighbor(indices,
data,
image_height,
image_width,
target_height,
target_width,
boxes=None,
box_indices=None,
extrapolation_value=None,
layout='NCHW',
coordinate_transformation_mode="align_corners",
out_dtype=None):
"""Perform resize operation with nearest neighbor method on the data.
For details about Nearest-neighbor interpolation please refer to
https://en.wikipedia.org/wiki/Nearest-neighbor_interpolation.
Parameters
----------
indices : tuple
The indices of input data
data : tvm.te.Tensor
inputs is a 4-D tensor with shape
[batch, channel, in_height, in_width]
or [batch, in_height, in_width, channel]
image_height : integer
Input image height
image_width : integer
Input image width
target_height : integer
The target resized image height
target_width : integer
The target resized image width
boxes : tvm.te.Tensor, optional
A 2-D tensor of shape [num_boxes, 4]. Each row of the tensor specifies
the coordinates of a box.
box_indices : tvm.te.Tensor, optional
A 1-D tensor of shape [num_boxes], box_indices[i] specifies the data that
the i-th box refers to.
extrapolation_value: float, optional
Value used for extrapolation, when applicable.
layout: string, optional
"NCHW", "NHWC", or "NCHWc".
coordinate_transformation_mode: string, optional
Describes how to transform the coordinate in the resized tensor
to the coordinate in the original tensor.
Refer to the ONNX Resize operator specification for details.
Available options are "half_pixel", "align_corners" and "asymmetric".
out_dtype: string, optional
Type to return. If left None will be same as input type.
Returns
-------
output : out_dtype
The computed result with type out_dtype
"""
def _cast_output(value, data_dtype="float32", out_dtype=None):
if out_dtype:
dtype = out_dtype
else:
dtype = data_dtype
return value.astype(dtype)
n, c, y, x, cc, inum, ic = get_2d_indices(indices, layout)
box_idx = box_indices(n) if box_indices is not None else n
if boxes is not None:
y1, x1 = boxes(n, 0), boxes(n, 1)
y2, x2 = boxes(n, 2), boxes(n, 3)
in_h = (image_height - 1) * (y2 - y1)
in_w = (image_width - 1) * (x2 - x1)
h_scale = in_h.astype('float') / (target_height - 1)
w_scale = in_w.astype('float') / (target_width - 1)
in_y = y1 * (image_height - 1) + h_scale * y
in_x = x1 * (image_width - 1) + w_scale * x
else:
if coordinate_transformation_mode == "align_corners":
h_scale = (image_height - 1).astype('float') / (target_height - 1)
w_scale = (image_width - 1).astype('float') / (target_width - 1)
elif coordinate_transformation_mode in ["asymmetric", "half_pixel"]:
h_scale = image_height.astype('float') / target_height
w_scale = image_width.astype('float') / target_width
else:
raise ValueError(
"Unsupported coordinate_transformation_mode: {}".format(
coordinate_transformation_mode))
in_y = h_scale * y
in_x = w_scale * x
if coordinate_transformation_mode == "align_corners" or boxes is not None:
closest_x_index = te.round(in_x).astype("int32")
closest_y_index = te.round(in_y).astype("int32")
else:
# Add epsilon to floor to prevent gpu rounding errors.
epsilon = 1e-5
closest_y_index = te.floor(in_y + epsilon).astype('int32')
closest_x_index = te.floor(in_x + epsilon).astype('int32')
value = get_2d_pixel(data, layout, boxes, image_height, image_width,
box_idx, c, closest_y_index, closest_x_index, cc,
inum, ic)
if extrapolation_value is not None:
out = tvm.tir.if_then_else(
in_y < 0, extrapolation_value,
tvm.tir.if_then_else(in_y > image_height - 1, extrapolation_value,
value))
# use extrapolation_value if in_x is out of boundary
value = tvm.tir.if_then_else(
in_x < 0, extrapolation_value,
tvm.tir.if_then_else(in_x > image_width - 1, extrapolation_value,
out))
return _cast_output(value, data.dtype, out_dtype=out_dtype)
def _nearest_sample(n, c, d, h, w):
z, y, x = _compute_source_index(n, d, h, w)
z_new = te.round(z).astype("int32")
y_new = te.round(y).astype("int32")
x_new = te.round(x).astype("int32")
return _get_pixel_value(n, c, z_new, y_new, x_new)
