diff --git a/include/tvm/relay/attrs/transform.h b/include/tvm/relay/attrs/transform.h index b5333961ebf9..eb22c70098c9 100644 --- a/include/tvm/relay/attrs/transform.h +++ b/include/tvm/relay/attrs/transform.h @@ -605,6 +605,17 @@ struct StftAttrs : public tvm::AttrsNode { } }; // struct StftAttrs +/*! \brief Attributes used in DFT operator */ +struct DFTAttrs : public tvm::AttrsNode { + Bool inverse = Bool(false); + + TVM_DECLARE_ATTRS(DFTAttrs, "relay.attrs.DFTAttrs") { + TVM_ATTR_FIELD(inverse) + .describe("Whether to perform the inverse discrete Fourier transform") + .set_default(Bool(false)); + } +}; // struct DFTAttrs + struct TriluAttrs : public tvm::AttrsNode { bool upper; diff --git a/python/tvm/relay/frontend/onnx.py b/python/tvm/relay/frontend/onnx.py index 2a1890627225..d3a542282c08 100644 --- a/python/tvm/relay/frontend/onnx.py +++ b/python/tvm/relay/frontend/onnx.py @@ -4877,6 +4877,116 @@ def _impl_v1(cls, inputs, attr, params): return mm_out +class DFT(OnnxOpConverter): + """Operator converter for discrete Fourier transform (DFT).""" + + @classmethod + def _impl_v17(cls, inputs, attr, params): + # ************************* Read attrs ************************* + axis = attr.get("axis") + inverse = attr.get("inverse") + onesided = attr.get("onesided") + + # ************************* Read inputs ************************ + input_tensor = inputs[0] + dft_length = inputs[1] + + # ************************* Parse inputs *********************** + t1 = ["float16", "float32", "float64"] + t2 = ["int32", "int64"] + + # input + assert infer_type(input_tensor).checked_type.dtype in t1 + input_shape = infer_shape(input_tensor) + assert len(input_shape) >= 3 + if axis < 0: + axis = len(input_shape) + axis + assert 1 <= axis <= len(input_shape) - 1, "axis is out of bounds" + + # dft_length + if dft_length is None: + dft_length = input_shape[axis] + else: + dft_length_dtype = infer_type(dft_length).checked_type.dtype + assert dft_length_dtype in t2 + dft_length = int(infer_value(dft_length, params).numpy()) + + # ************************ + input_tensor = cls._maybe_crop_or_pad(input_tensor, axis, dft_length) + + swap_axis = -1 + re_input_tensor, im_input_tensor = cls._split_real_and_imag_parts(input_tensor) + + re_input_tensor = cls._swap_axes(re_input_tensor, axis, swap_axis) + im_input_tensor = cls._swap_axes(im_input_tensor, axis, swap_axis) + + re_input_tensor, im_input_tensor = _op.dft(re_input_tensor, im_input_tensor, inverse) + + re_input_tensor = cls._swap_axes(re_input_tensor, axis, swap_axis) + im_input_tensor = cls._swap_axes(im_input_tensor, axis, swap_axis) + + if onesided: + re_input_tensor = cls._crop_onesided(re_input_tensor, axis) + im_input_tensor = cls._crop_onesided(im_input_tensor, axis) + + return cls._merge_real_and_imag_parts(re_input_tensor, im_input_tensor) + + @classmethod + def _crop_axis(cls, tensor, axis, new_dim): + shape = infer_shape(tensor) + slices = [slice(0, a, 1) for a in shape] + slices[axis] = slice(0, new_dim, 1) + return _op.strided_slice( + tensor, + begin=[s.start for s in slices], + end=[s.stop for s in slices], + strides=[s.step for s in slices], + axes=list(range(len(shape))), + ) + + @classmethod + def _maybe_crop_or_pad(cls, input_tensor, axis, n_fft): + shape = infer_shape(input_tensor) + if shape[axis] != n_fft: + if shape[axis] > n_fft: + return cls._crop_axis(input_tensor, axis, n_fft) + else: + pad_width = [(0, 0)] * len(shape) + pad_width[axis] = (0, n_fft - shape[axis]) + return _op.nn.pad(input_tensor, pad_width) + return input_tensor + + @classmethod + def _swap_axes(cls, tensor, axis1, axis2): + permutation = list(range(len(infer_shape(tensor)))) + permutation[axis1] = axis2 + permutation[axis2] = axis1 + return _op.transpose(tensor, permutation) + + @classmethod + def _split_real_and_imag_parts(cls, tensor): + shape = infer_shape(tensor) + dtype = infer_type(tensor).checked_type.dtype + if shape[-1] == 1: + re = tensor + im = _op.const(np.zeros(shape), dtype=dtype) + else: + re, im = _op.split(tensor, 2, -1) + + return _op.squeeze(re, -1), _op.squeeze(im, -1) + + @classmethod + def _merge_real_and_imag_parts(cls, re, im): + re = _op.expand_dims(re, axis=-1) + im = _op.expand_dims(im, axis=-1) + return _op.concatenate([re, im], axis=-1) + + @classmethod + def _crop_onesided(cls, tensor, axis): + shape = infer_shape(tensor) + return cls._crop_axis(tensor, axis, shape[axis] // 2 + 1) + + class NonMaxSuppression(OnnxOpConverter): """Operator converter for NonMaxSuppression.""" @@ -6696,6 +6806,7 @@ def _get_convert_map(opset): "Scan": Scan.get_converter(opset), # ML "LinearRegressor": LinearRegressor.get_converter(opset), + "DFT": DFT.get_converter(opset), # Sequence operators "SequenceConstruct": SequenceConstruct.get_converter(opset), "SequenceEmpty": SequenceEmpty.get_converter(opset), diff --git a/python/tvm/relay/op/_transform.py b/python/tvm/relay/op/_transform.py index e40179ed2d03..140b9835df6d 100644 --- a/python/tvm/relay/op/_transform.py +++ b/python/tvm/relay/op/_transform.py @@ -191,6 +191,20 @@ def stft_shape_func(attrs, inputs, _): ] +# DFT +@_reg.register_compute("dft") +def compute_dft(attrs, inputs, _): + """Compute definition of DFT""" + return topi.dft( + inputs[0], + inputs[1], + attrs.inverse, + ) + + +_reg.register_strategy("dft", strategy.dft_strategy) + + # trilu _reg.register_strategy("trilu", strategy.trilu_strategy) diff --git a/python/tvm/relay/op/strategy/cuda.py b/python/tvm/relay/op/strategy/cuda.py index fc1691fe9ef0..01c0654c3eb5 100644 --- a/python/tvm/relay/op/strategy/cuda.py +++ b/python/tvm/relay/op/strategy/cuda.py @@ -1412,3 +1412,14 @@ def stft_strategy_cuda(attrs, inputs, out_type, target): name="stft.cuda", ) return strategy + + +@dft_strategy.register(["cuda", "gpu"]) +def dft_strategy_cuda(attrs, inputs, out_type, target): + strategy = _op.OpStrategy() + strategy.add_implementation( + wrap_compute_dft(topi.cuda.dft), + wrap_topi_schedule(topi.generic.schedule_extern), + name="dft.cuda", + ) + return strategy diff --git a/python/tvm/relay/op/strategy/generic.py b/python/tvm/relay/op/strategy/generic.py index 4e0448f1799b..4811cae2ab7f 100644 --- a/python/tvm/relay/op/strategy/generic.py +++ b/python/tvm/relay/op/strategy/generic.py @@ -1468,6 +1468,32 @@ def _compute_stft(attrs, inputs, output_type): return _compute_stft +# dft +@override_native_generic_func("dft_strategy") +def dft_strategy(attrs, outs, out_type, target): + """DFT generic strategy""" + strategy = _op.OpStrategy() + strategy.add_implementation( + wrap_compute_dft(topi.dft), + wrap_topi_schedule(topi.generic.schedule_extern), + name="dft.generic", + ) + return strategy + + +def wrap_compute_dft(topi_compute): + """Wrap DFT compute""" + + def _compute_dft(attrs, inputs, _): + return topi_compute( + inputs[0], + inputs[1], + attrs.inverse, + ) + + return _compute_dft + + # trilu @override_native_generic_func("trilu_strategy") def trilu_strategy(attrs, outs, out_type, target): diff --git a/python/tvm/relay/op/transform.py b/python/tvm/relay/op/transform.py index 3df13da04426..1b0cc3588014 100644 --- a/python/tvm/relay/op/transform.py +++ b/python/tvm/relay/op/transform.py @@ -1987,6 +1987,33 @@ def stft( return _make.stft(data, n_fft, hop_length, win_length, window, normalized, onesided) +def dft(re_data, im_data, inverse=False): + """ + Computes the discrete Fourier transform of input (calculation along the last axis). + This gives frequency components of the signal as they change over time. + + Parameters + ---------- + re_data : relay.Expr + N-D tensor, real part of the input signal. + + im_data : relay.Expr + N-D tensor, imaginary part of the input signal. + If the signal is real, then the values of this tensor are zeros. + + inverse : bool + Whether to perform the inverse discrete fourier transform. + + Returns + ------- + re_output : relay.Expr + The Fourier Transform of the input (Real part). + im_output : relay.Expr + The Fourier Transform of the input (Imaginary part). + """ + return TupleWrapper(_make.dft(re_data, im_data, inverse), 2) + + def trilu(data, k, upper=True): """Given a 2-D matrix or batches of 2-D matrices, returns the upper or lower triangular part of the tensor. diff --git a/python/tvm/topi/__init__.py b/python/tvm/topi/__init__.py index 75867136e09e..3584191b86cc 100644 --- a/python/tvm/topi/__init__.py +++ b/python/tvm/topi/__init__.py @@ -47,7 +47,7 @@ from .einsum import * from .unique import * from .searchsorted import * -from .stft import * +from .signal import * from . import generic from . import nn from . import x86 diff --git a/python/tvm/topi/cuda/__init__.py b/python/tvm/topi/cuda/__init__.py index b746c95c0fc1..a6ced5bcf9bc 100644 --- a/python/tvm/topi/cuda/__init__.py +++ b/python/tvm/topi/cuda/__init__.py @@ -61,4 +61,4 @@ from .transform import * from .unique import * from .searchsorted import * -from .stft import * +from .signal import * diff --git a/python/tvm/topi/cuda/stft.py b/python/tvm/topi/cuda/signal.py similarity index 60% rename from python/tvm/topi/cuda/stft.py rename to python/tvm/topi/cuda/signal.py index 573c2ae39956..d08f41ab8912 100644 --- a/python/tvm/topi/cuda/stft.py +++ b/python/tvm/topi/cuda/signal.py @@ -133,3 +133,99 @@ def gen_ir( name="stft_cuda", tag="stft_cuda", ) + + +def dft( + re_data: te.Tensor, + im_data: te.Tensor, + inverse: tir.IntImm, +): + """ + Computes the discrete Fourier transform of input (calculation along the last axis). + This gives frequency components of the signal as they change over time. + + Parameters + ---------- + re_data : relay.Expr + N-D tensor, real part of the input signal. + + im_data : relay.Expr + N-D tensor, imaginary part of the input signal. + If the signal is real, then the values of this tensor are zeros. + + inverse : bool + Whether to perform the inverse discrete fourier transform. + + Returns + ------- + re_output : relay.Expr + The Fourier Transform of the input (Real part). + im_output : relay.Expr + The Fourier Transform of the input (Imaginary part). + """ + + def gen_ir( + re_data_buf, + im_data_buf, + re_output_buf, + im_output_buf, + ): + ib = tir.ir_builder.create() + re_data_ptr = ib.buffer_ptr(re_data_buf) + im_data_ptr = ib.buffer_ptr(im_data_buf) + re_output_ptr = ib.buffer_ptr(re_output_buf) + im_output_ptr = ib.buffer_ptr(im_output_buf) + + shape = re_data.shape + n_fft = shape[len(shape) - 1] + base_range = 1 + for i in range(len(shape) - 1): + base_range *= shape[i] + + sign = -1 if inverse else 1 + factor = 1.0 / n_fft if inverse else 1.0 + + max_threads = _get_max_threads(base_range) + with ib.new_scope(): + nthread_tx = max_threads + nthread_bx = ceil_div(base_range, max_threads) + tx = te.thread_axis("threadIdx.x") + bx = te.thread_axis("blockIdx.x") + ib.scope_attr(tx, "thread_extent", nthread_tx) + ib.scope_attr(bx, "thread_extent", nthread_bx) + + tid = bx * max_threads + tx + with ib.if_scope(tid < base_range): + base_idx = tid * n_fft + with ib.for_range(0, n_fft) as n: + n_idx = base_idx + n + re_output_ptr[n_idx] = tir.Cast(re_output_ptr.dtype, 0) + im_output_ptr[n_idx] = tir.Cast(im_output_ptr.dtype, 0) + _w = sign * -2 * pi * n / n_fft + with ib.for_range(0, n_fft) as k: + k_idx = base_idx + k + w = _w * k + cos_w = tir.Cast(re_output_ptr.dtype, tir.cos(w)) + sin_w = tir.Cast(re_output_ptr.dtype, tir.sin(w)) + re_output_ptr[n_idx] += ( + re_data_ptr[k_idx] * cos_w - im_data_ptr[k_idx] * sin_w + ) + im_output_ptr[n_idx] += ( + re_data_ptr[k_idx] * sin_w + im_data_ptr[k_idx] * cos_w + ) + + re_output_ptr[n_idx] *= tir.Cast(re_output_ptr.dtype, factor) + im_output_ptr[n_idx] *= tir.Cast(im_output_ptr.dtype, factor) + + return ib.get() + + output_shape = [re_data.shape] * 2 + + return te.extern( + shape=output_shape, + inputs=[re_data, im_data], + fcompute=lambda ins, outs: gen_ir(ins[0], ins[1], outs[0], outs[1]), + dtype=[re_data.dtype, im_data.dtype], + name="dft_cuda", + tag="dft_cuda", + ) diff --git a/python/tvm/topi/stft.py b/python/tvm/topi/signal.py similarity index 62% rename from python/tvm/topi/stft.py rename to python/tvm/topi/signal.py index b59c0245a052..64a804a851ab 100644 --- a/python/tvm/topi/stft.py +++ b/python/tvm/topi/signal.py @@ -123,3 +123,85 @@ def gen_ir( name="stft_cpu", tag="stft_cpu", ) + + +def dft( + re_data: te.Tensor, + im_data: te.Tensor, + inverse: tir.IntImm, +): + """ + Computes the discrete Fourier transform of input (calculation along the last axis). + This gives frequency components of the signal as they change over time. + + Parameters + ---------- + re_data : relay.Expr + N-D tensor, real part of the input signal. + + im_data : relay.Expr + N-D tensor, imaginary part of the input signal. + If the signal is real, then the values of this tensor are zeros. + + inverse : bool + Whether to perform the inverse discrete fourier transform. + + Returns + ------- + re_output : relay.Expr + The Fourier Transform of the input (Real part). + im_output : relay.Expr + The Fourier Transform of the input (Imaginary part). + """ + + def gen_ir( + re_data_buf, + im_data_buf, + re_output_buf, + im_output_buf, + ): + ib = tir.ir_builder.create() + re_data_ptr = ib.buffer_ptr(re_data_buf) + im_data_ptr = ib.buffer_ptr(im_data_buf) + re_output_ptr = ib.buffer_ptr(re_output_buf) + im_output_ptr = ib.buffer_ptr(im_output_buf) + + shape = re_data.shape + n_fft = shape[len(shape) - 1] + base_range = 1 + for i in range(len(shape) - 1): + base_range *= shape[i] + + sign = -1 if inverse else 1 + factor = 1.0 / n_fft if inverse else 1.0 + + with ib.for_range(0, base_range, kind="parallel") as i: + base_idx = i * n_fft + with ib.for_range(0, n_fft) as n: + n_idx = base_idx + n + re_output_ptr[n_idx] = tir.Cast(re_output_ptr.dtype, 0) + im_output_ptr[n_idx] = tir.Cast(im_output_ptr.dtype, 0) + _w = sign * -2 * pi * n / n_fft + with ib.for_range(0, n_fft) as k: + k_idx = base_idx + k + w = _w * k + cos_w = tir.Cast(re_output_ptr.dtype, tir.cos(w)) + sin_w = tir.Cast(re_output_ptr.dtype, tir.sin(w)) + re_output_ptr[n_idx] += re_data_ptr[k_idx] * cos_w - im_data_ptr[k_idx] * sin_w + im_output_ptr[n_idx] += re_data_ptr[k_idx] * sin_w + im_data_ptr[k_idx] * cos_w + + re_output_ptr[n_idx] *= tir.Cast(re_output_ptr.dtype, factor) + im_output_ptr[n_idx] *= tir.Cast(im_output_ptr.dtype, factor) + + return ib.get() + + output_shape = [re_data.shape] * 2 + + return te.extern( + shape=output_shape, + inputs=[re_data, im_data], + fcompute=lambda ins, outs: gen_ir(ins[0], ins[1], outs[0], outs[1]), + dtype=[re_data.dtype, im_data.dtype], + name="dft_cpu", + tag="dft_cpu", + ) diff --git a/src/relay/op/tensor/transform.cc b/src/relay/op/tensor/transform.cc index 01e5a7f5f359..806a17442903 100644 --- a/src/relay/op/tensor/transform.cc +++ b/src/relay/op/tensor/transform.cc @@ -1922,6 +1922,55 @@ RELAY_REGISTER_OP("stft") .set_support_level(3) .set_attr("TOpPattern", kOpaque); +// DFT +TVM_REGISTER_NODE_TYPE(DFTAttrs); +bool DFTRel(const Array& types, int num_inputs, const Attrs& attrs, + const TypeReporter& reporter) { + // types: [re_data, im_data, output] + ICHECK_EQ(types.size(), 3) + << "DFT: expects three types, two for the input and one for the output"; + ICHECK_EQ(num_inputs, 2) << "DFT: expect 2 inputs but " << num_inputs << " provided"; + const auto* re_data = types[0].as(); + const auto* im_data = types[1].as(); + + if (re_data == nullptr) { + ICHECK(types[0].as()) + << "DFT: expect re_data type to be TensorType but get " << types[0]; + return false; + } + if (im_data == nullptr) { + ICHECK(types[1].as()) + << "DFT: expect im_data type to be TensorType but get " << types[1]; + return false; + } + + std::vector shapes; + shapes.push_back(TensorType(re_data->shape, re_data->dtype)); + shapes.push_back(TensorType(im_data->shape, im_data->dtype)); + + reporter->Assign(types[2], TupleType(Array(shapes))); + + return true; +} + +Expr MakeDFT(Expr re_data, Expr im_data, Bool inverse) { + auto attrs = make_object(); + attrs->inverse = inverse; + static const Op& op = Op::Get("dft"); + return Call(op, {re_data, im_data}, Attrs(attrs), {}); +} + +TVM_REGISTER_GLOBAL("relay.op._make.dft").set_body_typed(MakeDFT); + +RELAY_REGISTER_OP("dft") + .describe(R"doc(Computes the discrete Fourier transform of input.)doc" TVM_ADD_FILELINE) + .set_num_inputs(2) + .add_argument("re_data", "Tensor", "Real part of input tensor.") + .add_argument("im_data", "Tensor", "Imaginary part of input tensor.") + .set_support_level(3) + .set_attr("TOpPattern", kOpaque) + .add_type_rel("DFT", DFTRel); + // meshgrid operator TVM_REGISTER_NODE_TYPE(MeshgridAttrs); diff --git a/tests/python/frontend/onnx/test_forward.py b/tests/python/frontend/onnx/test_forward.py index 293f4d38e649..116c023caadb 100644 --- a/tests/python/frontend/onnx/test_forward.py +++ b/tests/python/frontend/onnx/test_forward.py @@ -5429,9 +5429,6 @@ def verify_eyelike(indata, dynamic=False): "test_cumsum_2d_negative_axis", "test_det_2d", "test_det_nd", - "test_dft", - "test_dft_axis", - "test_dft_inverse", "test_dropout_default", "test_dropout_default_mask", "test_dropout_default_mask_ratio", @@ -5591,6 +5588,9 @@ def test_onnx_nodes(target, dev, onnx_test): # satisfies onnx precision for bicubic interpolation atol = 1e-4 + if "dft" in test_dir: + atol = 1e-3 + model = onnx.load(os.path.join(test_dir, "model.onnx")) for test_data_dir in glob.glob(os.path.join(test_dir, "test_data_set*")): inputs = [] @@ -7933,6 +7933,78 @@ def verify_linear_regressor(a_shape, c_shape, i_shape, targets=1, batch=1): verify_linear_regressor((1, 4), (3), (1)) +@tvm.testing.parametrize_targets +def test_dft(target, dev): + """test_dft""" + + def verify_dft( + _axis, + _inverse, + _onesided, + _dft_length, + _input_shape, + _output_shape, + ): + input_names = ["input"] + if _dft_length is not None: + input_names.append("dft_length") + + node = onnx.helper.make_node( + "DFT", + inputs=input_names, + outputs=["output"], + axis=_axis, + inverse=_inverse, + onesided=_onesided, + ) + + nodes = [] + if _dft_length is not None: + nodes.append( + make_constant_node("dft_length", TensorProto.INT32, [], [_dft_length]), + ) + nodes.append(node) + + graph = helper.make_graph( + nodes, + "dft_test", + inputs=[ + helper.make_tensor_value_info("input", TensorProto.FLOAT, _input_shape), + ], + outputs=[ + helper.make_tensor_value_info("output", TensorProto.FLOAT, _output_shape), + ], + ) + + model = helper.make_model(graph, producer_name="dft_test") + + _input = np.random.normal(size=_input_shape).astype("float32") + verify_with_ort_with_inputs( + model, + [_input], + [_input_shape], + target=target, + dev=dev, + rtol=1e-4, + atol=1e-4, + use_vm=False, + ) + + batch_size = 5 + n = 2 + D = 7 + + for axis in list(range(1, n)) + [-2]: + for inverse, onesided in [(0, 0), (0, 1), (1, 0)]: + for n_fft in [D, D - 1, D + 1]: + for c in [1, 2]: + input_shape = [batch_size] + n * [D] + [c] + output_shape = [batch_size] + n * [D] + [2] + if onesided == 1: + output_shape[axis] = output_shape[axis] // 2 + 1 + verify_dft(axis, inverse, onesided, n_fft, input_shape, output_shape) + + @tvm.testing.parametrize_targets def test_sequence(target, dev): """test_sequence""" diff --git a/tests/python/topi/python/test_topi_dft.py b/tests/python/topi/python/test_topi_dft.py new file mode 100644 index 000000000000..abab272e601d --- /dev/null +++ b/tests/python/topi/python/test_topi_dft.py @@ -0,0 +1,88 @@ +# 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. +"""Test code for discrete Fourier transform.""" +import numpy as np +import tvm +import tvm.testing +from tvm import topi +import tvm.topi.testing + + +inverse = tvm.testing.parameter(False, True) +shape = tvm.testing.parameter((7,), (3, 7), (3, 4, 5)) +dtype = tvm.testing.parameter("float16", "float32", "float64") + + +# pylint: disable=redefined-outer-name, invalid-name +def numpy_reference(inverse, re: np.ndarray, im: np.ndarray): + if inverse: + reference = np.fft.ifft(re + 1j * im) + else: + reference = np.fft.fft(re + 1j * im) + return np.real(reference), np.imag(reference) + + +def test_dft(target, dev, inverse, shape, dtype): + """Test for discrete Fourier transform.""" + implementations = { + "generic": ( + topi.dft, + topi.generic.schedule_extern, + ), + "gpu": ( + topi.cuda.dft, + topi.cuda.schedule_extern, + ), + "nvptx": ( + topi.cuda.dft, + topi.cuda.schedule_extern, + ), + } + + Re = tvm.te.placeholder(shape, dtype=dtype, name="Re") + Im = tvm.te.placeholder(shape, dtype=dtype, name="Im") + + with tvm.target.Target(target): + fcompute, fschedule = tvm.topi.testing.dispatch(target, implementations) + + outs = fcompute(Re, Im, inverse) + s = fschedule(outs) + + f = tvm.build(s, [Re, Im, *outs], target) + + re_np = np.random.normal(size=shape).astype(dtype) + im_np = np.random.normal(size=shape).astype(dtype) + + re = tvm.nd.array(re_np, device=dev) + im = tvm.nd.array(im_np, device=dev) + re_out = tvm.nd.array(np.zeros(shape).astype(dtype), device=dev) + im_out = tvm.nd.array(np.zeros(shape).astype(dtype), device=dev) + + f(re, im, re_out, im_out) + + re_reference, im_reference = numpy_reference(inverse, re_np, im_np) + + atol = rtol = 1e-3 + if dtype == "float16": + atol = rtol = 1e-1 + + tvm.testing.assert_allclose(re_out.numpy(), re_reference, rtol=rtol, atol=atol) + tvm.testing.assert_allclose(im_out.numpy(), im_reference, rtol=rtol, atol=atol) + + +if __name__ == "__main__": + tvm.testing.main()