From 28057b86600d323117d929ed21bd716ae557b79e Mon Sep 17 00:00:00 2001 From: Samuel Date: Mon, 11 May 2020 06:02:00 +0530 Subject: [PATCH] [TOPI][RELAY][TENSORFLOW]Math ops added (#5502) * [TOPI][RELAY][TENSORFLOW]Math ops added * Extra newline removed * CI fix * Review comments fixed * Review comments fixed --- python/tvm/relay/frontend/tensorflow.py | 7 ++ python/tvm/relay/op/_tensor.py | 7 +- python/tvm/relay/op/_tensor_grad.py | 53 +++++++++++-- python/tvm/relay/op/tensor.py | 75 ++++++++++++++++++ python/tvm/te/__init__.py | 3 +- src/relay/op/tensor/unary.cc | 55 ++++++++++++++ tests/python/frontend/tensorflow/test_forward.py | 97 +++++++----------------- tests/python/relay/test_op_grad_level1.py | 9 ++- topi/include/topi/elemwise.h | 5 ++ topi/python/topi/math.py | 84 ++++++++++++++++++++ topi/src/elemwise.cc | 25 ++++++ 11 files changed, 340 insertions(+), 80 deletions(-) diff --git a/python/tvm/relay/frontend/tensorflow.py b/python/tvm/relay/frontend/tensorflow.py index 120631e..913a016 100644 --- a/python/tvm/relay/frontend/tensorflow.py +++ b/python/tvm/relay/frontend/tensorflow.py @@ -1947,6 +1947,8 @@ _freezed_graph_pruned_op_list = ['ReadVariableOp', 'ResourceGather', 'Variable', # for N to 1 mapping, currently not supported(?) _convert_map = { 'Abs' : AttrCvt('abs'), + 'Acos' : AttrCvt('acos'), + 'Acosh' : AttrCvt('acosh'), 'Add' : _elemwise('add'), 'AddN' : _add_n(), 'AddV2' : _elemwise('add'), @@ -1954,8 +1956,11 @@ _convert_map = { 'Any' : _reduce('any'), 'ArgMax' : _argx(_op.argmax, 'argmax'), 'ArgMin' : _argx(_op.argmin, 'argmin'), + 'Asin' : AttrCvt('asin'), + 'Asinh' : AttrCvt('asinh'), 'Assert' : _assert(), 'Atan' : AttrCvt('atan'), + 'Atanh' : AttrCvt('atanh'), 'Atan2' : _atan2(), 'AvgPool' : _pooling('avg_pool'), 'AvgPool3D' : _pool3d('avg_pool3d'), @@ -1975,6 +1980,7 @@ _convert_map = { 'Conv2DBackpropInput' : _conv('conv_transpose'), 'Conv3D' : _conv3d('conv'), 'Cos' : AttrCvt('cos'), + 'Cosh' : AttrCvt('cosh'), 'CropAndResize' : _crop_and_resize(), 'DecodeJpeg' : _decode_image(), 'DepthToSpace' : _depth_to_space(), @@ -2051,6 +2057,7 @@ _convert_map = { 'Sigmoid' : AttrCvt('sigmoid'), 'Sign' : AttrCvt('sign'), 'Sin' : AttrCvt('sin'), + 'Sinh' : AttrCvt('sinh'), 'Size' : _size(), 'Slice' : _slice(), 'Softmax' : _softmax(), diff --git a/python/tvm/relay/op/_tensor.py b/python/tvm/relay/op/_tensor.py index ac006a4..e029e0c 100644 --- a/python/tvm/relay/op/_tensor.py +++ b/python/tvm/relay/op/_tensor.py @@ -16,10 +16,10 @@ # under the License. #pylint: disable=invalid-name, unused-argument, len-as-condition """Backend compiler related feature registration""" -import topi from tvm.runtime import convert from tvm.te.hybrid import script +import topi from topi.util import get_const_tuple from .op import register_compute, register_shape_func from .op import register_broadcast_schedule, register_injective_schedule @@ -34,7 +34,12 @@ register_broadcast_schedule("cos") register_broadcast_schedule("cosh") register_broadcast_schedule("sin") register_broadcast_schedule("sinh") +register_broadcast_schedule("acos") +register_broadcast_schedule("acosh") +register_broadcast_schedule("asin") +register_broadcast_schedule("asinh") register_broadcast_schedule("atan") +register_broadcast_schedule("atanh") register_broadcast_schedule("exp") register_broadcast_schedule("erf") register_broadcast_schedule("sqrt") diff --git a/python/tvm/relay/op/_tensor_grad.py b/python/tvm/relay/op/_tensor_grad.py index 07162c5..c034bcc 100644 --- a/python/tvm/relay/op/_tensor_grad.py +++ b/python/tvm/relay/op/_tensor_grad.py @@ -35,6 +35,7 @@ from .tensor import ( power, sin, sinh, + sqrt, zeros_like, equal, shape_of, @@ -98,10 +99,9 @@ def cos_grad(orig, grad): @register_gradient("cosh") def cosh_grad(orig, grad): - """Returns [grad * (-sinh(x))]""" + """Returns [grad * sinh(x)]""" x = orig.args[0] - ones = ones_like(x) - return [grad * (-ones * sinh(x))] + return [grad * sinh(x)] @register_gradient("sin") @@ -110,18 +110,61 @@ def sin_grad(orig, grad): x = orig.args[0] return [grad * cos(x)] + @register_gradient("sinh") def sinh_grad(orig, grad): """Returns [grad * cosh(x)]""" x = orig.args[0] return [grad * cosh(x)] + +@register_gradient("acos") +def acos_grad(orig, grad): + """Returns [grad * -1/((1 - (x ^ 2)) ^ 1/2)]""" + x = orig.args[0] + ones = ones_like(x) + return [grad * (-ones / sqrt(ones - (x * x)))] + + +@register_gradient("acosh") +def acosh_grad(orig, grad): + """Returns [grad * 1/((x - 1) ^ 1/2 * (x + 1) ^ 1/2)]""" + x = orig.args[0] + ones = ones_like(x) + return [grad * ones / sqrt((x * x) - ones)] + + +@register_gradient("asin") +def asin_grad(orig, grad): + """Returns [grad * 1/((1 - (x ^ 2)) ^ (1/2))]""" + x = orig.args[0] + ones = ones_like(x) + return [grad * ones / sqrt(ones - (x * x))] + + +@register_gradient("asinh") +def asinh_grad(orig, grad): + """Returns [grad * 1/((1 + (x ^ 2)) ^ (1/2))]""" + x = orig.args[0] + ones = ones_like(x) + return [grad * ones / sqrt(ones + (x * x))] + + @register_gradient("atan") def atan_grad(orig, grad): """Returns [grad * 1 / (1 + x ^ 2)]""" x = orig.args[0] - a = const(2.0) - return [grad * ones_like(x) / (ones_like(x) + power(x, a))] + ones = ones_like(x) + return [grad * ones / (ones + (x * x))] + + +@register_gradient("atanh") +def atanh_grad(orig, grad): + """Returns [grad * 1 / (1 - x ^ 2)]""" + x = orig.args[0] + ones = ones_like(x) + return [grad * ones / (ones - (x * x))] + @register_gradient("exp") def exp_grad(orig, grad): diff --git a/python/tvm/relay/op/tensor.py b/python/tvm/relay/op/tensor.py index a50a258..d5ae5cd 100644 --- a/python/tvm/relay/op/tensor.py +++ b/python/tvm/relay/op/tensor.py @@ -152,6 +152,66 @@ def sinh(data): """ return _make.sinh(data) +def acos(data): + """Compute elementwise acos of data. + + Parameters + ---------- + data : relay.Expr + The input data + + Returns + ------- + result : relay.Expr + The computed result. + """ + return _make.acos(data) + +def acosh(data): + """Compute elementwise acosh of data. + + Parameters + ---------- + data : relay.Expr + The input data + + Returns + ------- + result : relay.Expr + The computed result. + """ + return _make.acosh(data) + +def asin(data): + """Compute elementwise asin of data. + + Parameters + ---------- + data : relay.Expr + The input data + + Returns + ------- + result : relay.Expr + The computed result. + """ + return _make.asin(data) + +def asinh(data): + """Compute elementwise asinh of data. + + Parameters + ---------- + data : relay.Expr + The input data + + Returns + ------- + result : relay.Expr + The computed result. + """ + return _make.asinh(data) + def atan(data): """Compute elementwise atan of data. @@ -167,6 +227,21 @@ def atan(data): """ return _make.atan(data) +def atanh(data): + """Compute elementwise atanh of data. + + Parameters + ---------- + data : relay.Expr + The input data + + Returns + ------- + result : relay.Expr + The computed result. + """ + return _make.atanh(data) + def exp(data): """Compute elementwise exp of data. diff --git a/python/tvm/te/__init__.py b/python/tvm/te/__init__.py index d22b350..939956c 100644 --- a/python/tvm/te/__init__.py +++ b/python/tvm/te/__init__.py @@ -19,8 +19,9 @@ """ # expose all operators in tvm tir.op from tvm.tir import any, all, min_value, max_value, trace -from tvm.tir import exp, erf, tanh, sigmoid, log, tan, cos, sin, atan, sqrt, rsqrt, floor, ceil +from tvm.tir import exp, erf, tanh, sigmoid, log, tan, cos, sin, sqrt, rsqrt, floor, ceil from tvm.tir import sinh, cosh, log2, log10 +from tvm.tir import asin, asinh, acos, acosh, atan, atanh from tvm.tir import trunc, abs, round, nearbyint, power, popcount, fmod, if_then_else from tvm.tir import isnan, isfinite, isinf from tvm.tir import div, indexdiv, indexmod, truncdiv, truncmod, floordiv, floormod diff --git a/src/relay/op/tensor/unary.cc b/src/relay/op/tensor/unary.cc index 2d397ab..5714379 100644 --- a/src/relay/op/tensor/unary.cc +++ b/src/relay/op/tensor/unary.cc @@ -128,6 +128,50 @@ RELAY_REGISTER_UNARY_OP("sinh") .set_attr("FTVMCompute", RELAY_UNARY_COMPUTE(topi::sinh)); +RELAY_REGISTER_UNARY_OP("acos") +.describe(R"code(Returns the acos of input array, computed element-wise. + +.. math:: + Y = acos(X) + +)code" TVM_ADD_FILELINE) +.set_support_level(1) +.set_attr("FTVMCompute", RELAY_UNARY_COMPUTE(topi::acos)); + + +RELAY_REGISTER_UNARY_OP("acosh") +.describe(R"code(Returns the acosh of input array, computed element-wise. + +.. math:: + Y = acosh(X) + +)code" TVM_ADD_FILELINE) +.set_support_level(1) +.set_attr("FTVMCompute", RELAY_UNARY_COMPUTE(topi::acosh)); + + +RELAY_REGISTER_UNARY_OP("asin") +.describe(R"code(Returns the asin of input array, computed element-wise. + +.. math:: + Y = asin(X) + +)code" TVM_ADD_FILELINE) +.set_support_level(1) +.set_attr("FTVMCompute", RELAY_UNARY_COMPUTE(topi::asin)); + + +RELAY_REGISTER_UNARY_OP("asinh") +.describe(R"code(Returns the asinh of input array, computed element-wise. + +.. math:: + Y = asinh(X) + +)code" TVM_ADD_FILELINE) +.set_support_level(1) +.set_attr("FTVMCompute", RELAY_UNARY_COMPUTE(topi::asinh)); + + RELAY_REGISTER_UNARY_OP("atan") .describe(R"code(Returns the atan of input array, computed element-wise. @@ -139,6 +183,17 @@ RELAY_REGISTER_UNARY_OP("atan") .set_attr("FTVMCompute", RELAY_UNARY_COMPUTE(topi::atan)); +RELAY_REGISTER_UNARY_OP("atanh") +.describe(R"code(Returns the atanh of input array, computed element-wise. + +.. math:: + Y = atanh(X) + +)code" TVM_ADD_FILELINE) +.set_support_level(1) +.set_attr("FTVMCompute", RELAY_UNARY_COMPUTE(topi::atanh)); + + RELAY_REGISTER_UNARY_OP("exp") .describe(R"code(Returns the exp input array, computed element-wise. diff --git a/tests/python/frontend/tensorflow/test_forward.py b/tests/python/frontend/tensorflow/test_forward.py index 858ef6b..cd6c454 100644 --- a/tests/python/frontend/tensorflow/test_forward.py +++ b/tests/python/frontend/tensorflow/test_forward.py @@ -2595,15 +2595,6 @@ def test_forward_zeros_like(): _test_forward_zeros_like((2, 3, 11), "float64") -def test_forward_erf(): - ishape = (1, 3, 10, 10) - inp_array = np.random.uniform(-5, 5, size=ishape).astype(np.float32) - with tf.Graph().as_default(): - in1 = tf.placeholder(shape=inp_array.shape, dtype=inp_array.dtype) - tf.math.erf(in1) - compare_tf_with_tvm(inp_array, 'Placeholder:0', 'Erf:0') - - def test_forward_squared_difference(): ishape = (1, 3, 10, 14) inp_array_a = np.random.uniform(-5, 5, size=ishape).astype(np.float32) @@ -2670,52 +2661,32 @@ def test_forward_pow_exp(): compare_tf_with_tvm([np_in1], ['in1:0'], 'exp:0') -def test_forward_log(): - """test operator Log """ - np_data = np.random.uniform(1, 100, size=(2, 3, 5)).astype(np.float32) - tf.reset_default_graph() - with tf.Graph().as_default(): - in_data = tf.placeholder(tf.float32, (2, 3, 5), name="in_data") - tf.log(in_data, name="log") - compare_tf_with_tvm([np_data], ['in_data:0'], 'log:0') - - -def test_forward_log1p(): - """test operator Log1p """ - np_data = np.random.uniform(1, 100, size=(2, 3, 5)).astype(np.float32) - tf.reset_default_graph() - with tf.Graph().as_default(): - in_data = tf.placeholder(tf.float32, (2, 3, 5), name="in_data") - tf.log1p(in_data, name="log1p") - compare_tf_with_tvm([np_data], ['in_data:0'], 'log1p:0') - - -def test_forward_cos(): - """test operator cos """ - np_data = np.random.uniform(1, 100, size=(2, 3, 5)).astype(np.float32) - tf.reset_default_graph() - with tf.Graph().as_default(): - in_data = tf.placeholder(tf.float32, (2, 3, 5), name="in_data") - tf.cos(in_data, name="cos") - compare_tf_with_tvm([np_data], ['in_data:0'], 'cos:0') - +def test_forward_unary(): + def _test_forward_unary(op, a_min=1, a_max=5, dtype=np.float32): + """test unary operators""" + np_data = np.random.uniform(a_min, a_max, size=(2, 3, 5)).astype(dtype) + tf.reset_default_graph() + with tf.Graph().as_default(): + in_data = tf.placeholder(dtype, (2, 3, 5), name="in_data") + out = op(in_data) + compare_tf_with_tvm([np_data], ['in_data:0'], out.name) + + _test_forward_unary(tf.acos, -1, 1) + _test_forward_unary(tf.asin, -1, 1) + _test_forward_unary(tf.atanh, -1, 1) + _test_forward_unary(tf.sinh) + _test_forward_unary(tf.cosh) + _test_forward_unary(tf.acosh) + _test_forward_unary(tf.asinh) + _test_forward_unary(tf.atan) + _test_forward_unary(tf.sin) + _test_forward_unary(tf.cos) + _test_forward_unary(tf.tan) + _test_forward_unary(tf.tanh) + _test_forward_unary(tf.erf) + _test_forward_unary(tf.log) + _test_forward_unary(tf.log1p) -def test_forward_tan(): - """test operator tan """ - np_data = np.random.uniform(1, 100, size=(2, 3, 5)).astype(np.float32) - tf.reset_default_graph() - in_data = tf.placeholder(tf.float32, (2, 3, 5), name="in_data") - tf.tan(in_data, name="tan") - compare_tf_with_tvm([np_data], ['in_data:0'], 'tan:0') - -def test_forward_atan(): - """test operator tan """ - tf.disable_eager_execution() - np_data = np.random.uniform(1, 100, size=(2, 3, 5)).astype(np.float32) - tf.reset_default_graph() - in_data = tf.placeholder(tf.float32, (2, 3, 5), name="in_data") - tf.atan(in_data, name="atan") - compare_tf_with_tvm([np_data], ['in_data:0'], 'atan:0') def test_forward_atan2(): """test operator tan """ @@ -2729,16 +2700,6 @@ def test_forward_atan2(): compare_tf_with_tvm([np_data_1, np_data_2], ['in_data_1:0', 'in_data_2:0'], 'atan2:0') -def test_forward_sin(): - """test operator sin """ - np_data = np.random.uniform(1, 100, size=(2, 3, 5)).astype(np.float32) - tf.reset_default_graph() - with tf.Graph().as_default(): - in_data = tf.placeholder(tf.float32, (2, 3, 5), name="in_data") - tf.sin(in_data, name="sin") - compare_tf_with_tvm([np_data], ['in_data:0'], 'sin:0') - - def test_forward_negative(): """test tf operator Neg """ np_data = np.random.uniform(-100, 255, @@ -3238,7 +3199,6 @@ if __name__ == '__main__': test_forward_left_shift() test_forward_truncatemod() test_forward_one_hot() - test_forward_atan() test_forward_atan2() # Activations @@ -3254,11 +3214,6 @@ if __name__ == '__main__': test_forward_reverse_v2() test_forward_pow_exp() test_forward_sign() - test_forward_log() - test_forward_log1p() - test_forward_tan() - test_forward_cos() - test_forward_sin() test_forward_negative() test_forward_divide() test_forward_abs() @@ -3271,13 +3226,13 @@ if __name__ == '__main__': test_forward_log_softmax() test_forward_bias_add() test_forward_zeros_like() - test_forward_erf() test_forward_squared_difference() test_forward_add_n() test_forward_floormod() test_forward_isfinite() test_forward_isinf() test_forward_unravel_index() + test_forward_unary() # Reductions test_forward_argminmax() diff --git a/tests/python/relay/test_op_grad_level1.py b/tests/python/relay/test_op_grad_level1.py index 5ef98a5..9faf6d9 100644 --- a/tests/python/relay/test_op_grad_level1.py +++ b/tests/python/relay/test_op_grad_level1.py @@ -68,8 +68,13 @@ def test_unary_op(): (tvm.relay.atan, lambda x: 1 / (1 + np.power(x, 2.0))), (tvm.relay.log2, lambda x: 1 / (np.log(2) * x)), (tvm.relay.log10, lambda x: 1 / (np.log(10) * x)), - (tvm.relay.cosh, lambda x: -1.0 * np.sinh(x)), - (tvm.relay.sinh, lambda x: np.cosh(x))]: + (tvm.relay.cosh, lambda x: np.sinh(x)), + (tvm.relay.sinh, lambda x: np.cosh(x)), + (tvm.relay.asin, lambda x: 1. / (1. - x**2) ** (1./2.)), + (tvm.relay.acos, lambda x: -1. / (1. - x**2.) ** (1./2.)), + (tvm.relay.acosh, lambda x: 1./ (x**2 - 1.)**(1./2.)), + (tvm.relay.asinh, lambda x: 1./ (x**2 + 1.)**(1./2.)), + (tvm.relay.atanh, lambda x: -1./ (x**2 - 1.))]: check_single_op(opfunc, ref) diff --git a/topi/include/topi/elemwise.h b/topi/include/topi/elemwise.h index 11eda86..dfcf83f 100644 --- a/topi/include/topi/elemwise.h +++ b/topi/include/topi/elemwise.h @@ -61,7 +61,12 @@ TOPI_DECLARE_UNARY_OP(cosh); TOPI_DECLARE_UNARY_OP(tan); TOPI_DECLARE_UNARY_OP(sin); TOPI_DECLARE_UNARY_OP(sinh); +TOPI_DECLARE_UNARY_OP(acos); +TOPI_DECLARE_UNARY_OP(acosh); +TOPI_DECLARE_UNARY_OP(asin); +TOPI_DECLARE_UNARY_OP(asinh); TOPI_DECLARE_UNARY_OP(atan); +TOPI_DECLARE_UNARY_OP(atanh); TOPI_DECLARE_UNARY_OP(isnan); TOPI_DECLARE_UNARY_OP(tanh); TOPI_DECLARE_UNARY_OP(isfinite); diff --git a/topi/python/topi/math.py b/topi/python/topi/math.py index c95cbf8..d715308 100644 --- a/topi/python/topi/math.py +++ b/topi/python/topi/math.py @@ -195,6 +195,74 @@ def sinh(x): @tvm.te.tag_scope(tag=tag.ELEMWISE) +def acos(x): + """Take arc cos of input x. + + Parameters + ---------- + x : tvm.te.Tensor + Input argument. + + Returns + ------- + y : tvm.te.Tensor + The result. + """ + return te.compute(x.shape, lambda *i: te.acos(x(*i))) + + +@tvm.te.tag_scope(tag=tag.ELEMWISE) +def acosh(x): + """Take arc cosh of input x. + + Parameters + ---------- + x : tvm.te.Tensor + Input argument. + + Returns + ------- + y : tvm.te.Tensor + The result. + """ + return te.compute(x.shape, lambda *i: te.acosh(x(*i))) + + +@tvm.te.tag_scope(tag=tag.ELEMWISE) +def asin(x): + """Take arc sin of input x. + + Parameters + ---------- + x : tvm.te.Tensor + Input argument. + + Returns + ------- + y : tvm.te.Tensor + The result. + """ + return te.compute(x.shape, lambda *i: te.asin(x(*i))) + + +@tvm.te.tag_scope(tag=tag.ELEMWISE) +def asinh(x): + """Take arc sinh of input x. + + Parameters + ---------- + x : tvm.te.Tensor + Input argument. + + Returns + ------- + y : tvm.te.Tensor + The result. + """ + return te.compute(x.shape, lambda *i: te.asinh(x(*i))) + + +@tvm.te.tag_scope(tag=tag.ELEMWISE) def atan(x): """Take atan of input x. @@ -211,6 +279,22 @@ def atan(x): return te.compute(x.shape, lambda *i: te.atan(x(*i))) @tvm.te.tag_scope(tag=tag.ELEMWISE) +def atanh(x): + """Take atanh of input x. + + Parameters + ---------- + x : tvm.te.Tensor + Input argument. + + Returns + ------- + y : tvm.te.Tensor + The result. + """ + return te.compute(x.shape, lambda *i: te.atanh(x(*i))) + +@tvm.te.tag_scope(tag=tag.ELEMWISE) def floor(x): """Take floor of input x. diff --git a/topi/src/elemwise.cc b/topi/src/elemwise.cc index a19467c..2c59994 100644 --- a/topi/src/elemwise.cc +++ b/topi/src/elemwise.cc @@ -31,6 +31,31 @@ namespace topi { using namespace tvm; using namespace tvm::runtime; +TVM_REGISTER_GLOBAL("topi.acos") +.set_body([](TVMArgs args, TVMRetValue *rv) { + *rv = acos(args[0]); + }); + +TVM_REGISTER_GLOBAL("topi.acosh") +.set_body([](TVMArgs args, TVMRetValue *rv) { + *rv = acosh(args[0]); + }); + +TVM_REGISTER_GLOBAL("topi.asin") +.set_body([](TVMArgs args, TVMRetValue *rv) { + *rv = asin(args[0]); + }); + +TVM_REGISTER_GLOBAL("topi.asinh") +.set_body([](TVMArgs args, TVMRetValue *rv) { + *rv = asinh(args[0]); + }); + +TVM_REGISTER_GLOBAL("topi.atanh") +.set_body([](TVMArgs args, TVMRetValue *rv) { + *rv = atanh(args[0]); + }); + TVM_REGISTER_GLOBAL("topi.exp") .set_body([](TVMArgs args, TVMRetValue *rv) { *rv = exp(args[0]); -- 2.7.4