From 7e086b26a175e683ab074a76d09ca4aef2521577 Mon Sep 17 00:00:00 2001 From: tom__chen Date: Mon, 19 Apr 2021 15:58:51 -0400 Subject: [PATCH] resize_bilinear gpu op --- .../gpu/cuda_impl/resize_bilinear_impl.cu | 66 ++ .../gpu/cuda_impl/resize_bilinear_impl.cuh | 24 + .../gpu/nn/resize_bilinear_gpu_kernel.cc | 26 + .../gpu/nn/resize_bilinear_gpu_kernel.h | 128 ++++ mindspore/ops/operations/nn_ops.py | 2 +- tests/st/ops/gpu/test_resize_bilinear_op.py | 596 ++++++++++++++++++ 6 files changed, 841 insertions(+), 1 deletion(-) create mode 100644 mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cu create mode 100644 mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cuh create mode 100644 mindspore/ccsrc/backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.cc create mode 100644 mindspore/ccsrc/backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.h create mode 100644 tests/st/ops/gpu/test_resize_bilinear_op.py diff --git a/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cu b/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cu new file mode 100644 index 0000000000..ed716af641 --- /dev/null +++ b/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cu @@ -0,0 +1,66 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed 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. + */ + +#include "backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cuh" +#include "runtime/device/gpu/cuda_common.h" +#include "include/cuda_fp16.h" +template +__global__ void ResizeBilinear(const T *input, const int n, const int c, const int input_h, const int input_w, + const int output_h, const int output_w, const int nchw, const int chw, const int hw, const float h_scale, + const float w_scale, float *output) { + for (size_t pos = blockIdx.x * blockDim.x + threadIdx.x; pos < nchw; pos += blockDim.x * gridDim.x) { + const int posn = pos / chw; + const int posc = pos / hw % c; + const int posh = pos / output_w % output_h; + const int posw = pos % output_w; + const float posw_scaled = w_scale * posw; + const float posh_scaled = h_scale * posh; + const int w_low = max(static_cast(floorf(posw_scaled)), 0); // NOLINT + const int w_high = min(static_cast(ceilf(posw_scaled)), input_w - 1); // NOLINT + const int h_low = max(static_cast(floorf(posh_scaled)), 0); // NOLINT + const int h_high = min(static_cast(ceilf(posh_scaled)), input_h - 1); // NOLINT + const float w_alpha = posw_scaled - w_low; + const float w_beta = 1.0f - w_alpha; + const float h_alpha = posh_scaled - h_low; + const float h_beta = 1.0f - h_alpha; + const int input_start = input_h * input_w * (posn * c + posc); + const float p1 = static_cast(input[input_start + (h_low * input_w) + w_low]); + const float p2 = static_cast(input[input_start + (h_low * input_w) + w_high]); + const float p3 = static_cast(input[input_start + (h_high * input_w) + w_low]); + const float p4 = static_cast(input[input_start + (h_high * input_w) + w_high]); + output[pos] = (p1 * h_beta * w_beta) + (p2 * h_beta * w_alpha) + (p3 * h_alpha * w_beta) + (p4 * h_alpha * w_alpha); + } + return; +} + +template +void CalResizeBilinear(const T *input, const int n, const int c, const int input_h, const int input_w, + const int output_h, const int output_w, const float h_scale, const float w_scale, float *output, + cudaStream_t cuda_stream) { + const int nchw = n * c * output_h * output_w; + const int chw = c * output_h * output_w; + const int hw = output_h * output_w; + ResizeBilinear<<>>(input, n, c, input_h, input_w, output_h, + output_w, nchw, chw, hw, h_scale, w_scale, output); + return; +} + +template void CalResizeBilinear(const float *input, const int n, const int c, const int input_h, + const int input_w, const int output_h, const int output_w, const float h_scale, const float w_scale, float *output, + cudaStream_t cuda_stream); +template void CalResizeBilinear(const half *input, const int n, const int c, const int input_h, + const int input_w, const int output_h, const int output_w, const float h_scale, const float w_scale, float *output, + cudaStream_t cuda_stream); diff --git a/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cuh b/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cuh new file mode 100644 index 0000000000..d80405af3c --- /dev/null +++ b/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cuh @@ -0,0 +1,24 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed 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. + */ + +#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_RESIZE_BILINEAR_H_ +#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_RESIZE_BILINEAR_H_ +#include "runtime/device/gpu/cuda_common.h" +template +void CalResizeBilinear(const T *input, const int n_, const int c_, const int input_h_, const int input_w_, + const int output_h_, const int output_w_, const float h_scale, const float w_scale, float *output, + cudaStream_t cuda_stream); +#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_RESIZE_BILINEAR_H_ diff --git a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.cc b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.cc new file mode 100644 index 0000000000..8f2eab7e08 --- /dev/null +++ b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.cc @@ -0,0 +1,26 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed 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. + */ + +#include "backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.h" + +namespace mindspore { +namespace kernel { +MS_REG_GPU_KERNEL_ONE(ResizeBilinear, KernelAttr().AddInputAttr(kNumberTypeFloat32).AddOutputAttr(kNumberTypeFloat32), + ResizeBilinearGpuKernel, float) +MS_REG_GPU_KERNEL_ONE(ResizeBilinear, KernelAttr().AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat32), + ResizeBilinearGpuKernel, half) +} // namespace kernel +} // namespace mindspore diff --git a/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.h b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.h new file mode 100644 index 0000000000..958d07027f --- /dev/null +++ b/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/resize_bilinear_gpu_kernel.h @@ -0,0 +1,128 @@ +/** + * Copyright 2021 Huawei Technologies Co., Ltd + * + * Licensed 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. + */ + +#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_RESIZE_BILINEAR_GPU_KERNEL_H_ +#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_RESIZE_BILINEAR_GPU_KERNEL_H_ + +#include +#include "backend/kernel_compiler/gpu/gpu_kernel.h" +#include "backend/kernel_compiler/gpu/gpu_kernel_factory.h" +#include "backend/kernel_compiler/gpu/cuda_impl/resize_bilinear_impl.cuh" + +namespace mindspore { +namespace kernel { +template +class ResizeBilinearGpuKernel : public GpuKernel { + public: + ResizeBilinearGpuKernel() { ResetResource(); } + ~ResizeBilinearGpuKernel() override = default; + + const std::vector &GetInputSizeList() const override { return input_size_list_; } + const std::vector &GetOutputSizeList() const override { return output_size_list_; } + const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } + + bool Launch(const std::vector &inputs, const std::vector &workspace, + const std::vector &outputs, void *stream_ptr) override { + T *input = GetDeviceAddress(inputs, 0); + float *output = GetDeviceAddress(outputs, 0); + float h_scale = Scaling(input_h_, output_h_, align_corners_); + float w_scale = Scaling(input_w_, output_w_, align_corners_); + CalResizeBilinear(input, n_, c_, input_h_, input_w_, output_h_, output_w_, h_scale, w_scale, output, + reinterpret_cast(stream_ptr)); + return true; + } + + bool Init(const CNodePtr &kernel_node) override { + size_t input_num = AnfAlgo::GetInputTensorNum(kernel_node); + if (input_num != 1) { + MS_LOG(ERROR) << "Input number is " << input_num << ", but ResizeBilinear needs 1 input."; + return false; + } + size_t output_num = AnfAlgo::GetOutputTensorNum(kernel_node); + if (output_num != 1) { + MS_LOG(ERROR) << "Output number is " << output_num << ", but ResizeBilinear has 1 output."; + return false; + } + std::vector input_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0); + std::vector output_shape = AnfAlgo::GetOutputInferShape(kernel_node, 0); + if (input_shape.size() != 4) { + MS_LOG(ERROR) << "Input is " << input_shape.size() << "-D, but ResizeBilinear supports only 4-D inputs."; + return false; + } + n_ = SizeToInt(input_shape[0]); + c_ = SizeToInt(input_shape[1]); + input_h_ = SizeToInt(input_shape[2]); + input_w_ = SizeToInt(input_shape[3]); + output_h_ = SizeToInt(output_shape[2]); + output_w_ = SizeToInt(output_shape[3]); + input_size_ = sizeof(T); + for (auto x : input_shape) { + input_size_ *= x; + } + output_size_ = sizeof(float); + for (auto x : output_shape) { + output_size_ *= x; + } + align_corners_ = GetAttr(kernel_node, "align_corners"); + InitSizeLists(); + return true; + } + + void ResetResource() noexcept override { + align_corners_ = false; + n_ = 0; + c_ = 0; + input_h_ = 0; + input_w_ = 0; + output_h_ = 0; + output_w_ = 0; + input_size_ = 0; + output_size_ = 0; + workspace_size_ = 0; + input_size_list_.clear(); + output_size_list_.clear(); + workspace_size_list_.clear(); + } + + protected: + void InitSizeLists() override { + input_size_list_.push_back(input_size_); + output_size_list_.push_back(output_size_); + } + + private: + float Scaling(const int in_size, const int out_size, bool align_corners) { + return (align_corners && out_size > 1) ? (in_size - 1) / static_cast(out_size - 1) + : in_size / static_cast(out_size); + } + + bool align_corners_; + int n_; + int c_; + int input_h_; + int input_w_; + int output_h_; + int output_w_; + size_t input_size_; + size_t output_size_; + size_t workspace_size_; + std::vector input_size_list_; + std::vector output_size_list_; + std::vector workspace_size_list_; +}; +} // namespace kernel +} // namespace mindspore +#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_RESIZE_BILINEAR_GPU_KERNEL_H_ diff --git a/mindspore/ops/operations/nn_ops.py b/mindspore/ops/operations/nn_ops.py index 46d776ec97..82c89ef308 100644 --- a/mindspore/ops/operations/nn_ops.py +++ b/mindspore/ops/operations/nn_ops.py @@ -3194,7 +3194,7 @@ class ResizeBilinear(PrimitiveWithInfer): ValueError: If length of shape of `input` is not equal to 4. Supported Platforms: - ``Ascend`` ``CPU`` + ``Ascend`` ``CPU`` ``GPU`` Examples: >>> tensor = Tensor([[[[1, 2, 3, 4, 5], [1, 2, 3, 4, 5]]]], mindspore.float32) diff --git a/tests/st/ops/gpu/test_resize_bilinear_op.py b/tests/st/ops/gpu/test_resize_bilinear_op.py new file mode 100644 index 0000000000..dd4c54474a --- /dev/null +++ b/tests/st/ops/gpu/test_resize_bilinear_op.py @@ -0,0 +1,596 @@ +# Copyright 2021 Huawei Technologies Co., Ltd +# +# Licensed 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. +# ============================================================================ +import numpy as np +import pytest + +from mindspore import context, Tensor +from mindspore.ops import operations as P +from mindspore import nn + +context.set_context(mode=context.GRAPH_MODE, device_target="GPU") + + +class NetResizeBilinear(nn.Cell): + def __init__(self, size=None, align_corner=False): + super(NetResizeBilinear, self).__init__() + self.op = P.ResizeBilinear(size=size, align_corners=align_corner) + + def construct(self, inputs): + return self.op(inputs) + + +@pytest.mark.level0 +@pytest.mark.platform_x86_gpu_training +@pytest.mark.env_onecard +def test_resize_nn_grayscale_integer_ratio_half(datatype=np.float16): + input_tensor = Tensor(np.array( + [[[[0.1, 0.2, 0.3], [0.4, 0.5, 0.6], [0.7, 0.8, 0.9]]]]).astype(datatype)) + + # larger h and w + resize_nn = NetResizeBilinear((9, 9)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.13330078, 0.16662598, 0.19995117, 0.23331706, + 0.26668295, 0.30004883, 0.30004883, 0.30004883], + [0.19995117, 0.23328993, 0.26662868, 0.29996747, 0.33333334, + 0.36669925, 0.40006512, 0.40006512, 0.40006512], + [0.29992676, 0.33327907, 0.36663142, 0.39998373, 0.4333496, + 0.4667155, 0.5000814, 0.5000814, 0.5000814], + [0.39990234, 0.43326822, 0.46663412, 0.5, 0.5333659, + 0.5667318, 0.60009766, 0.60009766, 0.60009766], + [0.5, 0.5333116, 0.5666233, 0.59993494, 0.6333008, + 0.66666675, 0.7000326, 0.7000326, 0.7000326], + [0.60009766, 0.633355, 0.66661245, 0.6998698, 0.7332357, + 0.7666016, 0.79996747, 0.79996747, 0.79996747], + [0.7001953, 0.73339844, 0.76660156, 0.7998047, 0.8331706, + 0.8665365, 0.89990234, 0.89990234, 0.89990234], + [0.7001953, 0.73339844, 0.76660156, 0.7998047, 0.8331706, + 0.8665365, 0.89990234, 0.89990234, 0.89990234], + [0.7001953, 0.73339844, 0.76660156, 0.7998047, 0.8331706, + 0.8665365, 0.89990234, 0.89990234, 0.89990234]]]]).astype(np.float32)) + error = np.ones(shape=[9, 9]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h and w + resize_nn = NetResizeBilinear((1, 1)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559]]]]).astype(np.float32)) + error = np.ones(shape=[1, 1]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h, larger w + resize_nn = NetResizeBilinear((1, 6)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.09997559, 0.14996338, 0.19995117, 0.25, 0.30004883, 0.30004883]]]]).astype(np.float32)) + error = np.ones(shape=[1, 6]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # larger h, smaller w + resize_nn = NetResizeBilinear((6, 1)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.09997559], [0.24993896], [0.39990234], [0.5500488], [0.7001953], [0.7001953]]]]).astype( + np.float32)) + error = np.ones(shape=[6, 1]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h, same w + resize_nn = NetResizeBilinear((1, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.09997559, 0.19995117, 0.30004883]]]]).astype(np.float32)) + error = np.ones(shape=[1, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # larger h, same w + resize_nn = NetResizeBilinear((6, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.19995117, 0.30004883], + [0.24993896, 0.3499756, 0.45007324], + [0.39990234, 0.5, 0.60009766], + [0.5500488, 0.64990234, 0.75], + [0.7001953, 0.7998047, 0.89990234], + [0.7001953, 0.7998047, 0.89990234]]]]).astype(np.float32)) + error = np.ones(shape=[6, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same h, smaller w + resize_nn = NetResizeBilinear((3, 1)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.09997559], [0.39990234], [0.7001953]]]]).astype(np.float32)) + error = np.ones(shape=[3, 1]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same h, larger w + resize_nn = NetResizeBilinear((3, 6)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.14996338, 0.19995117, 0.25, 0.30004883, + 0.30004883], + [0.39990234, 0.44995117, 0.5, 0.5500488, 0.60009766, + 0.60009766], + [0.7001953, 0.75, 0.7998047, 0.8498535, 0.89990234, + 0.89990234]]]]).astype(np.float32)) + error = np.ones(shape=[3, 6]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same w, same h (identity) + resize_nn = NetResizeBilinear((3, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array( + [[[[0.09997559, 0.19995117, 0.30004883], + [0.39990234, 0.5, 0.60009766], + [0.7001953, 0.7998047, 0.89990234]]]]).astype(np.float32)) + error = np.ones(shape=[3, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + +@pytest.mark.level0 +@pytest.mark.platform_x86_gpu_training +@pytest.mark.env_onecard +def test_resize_nn_grayscale_integer_ratio_float(datatype=np.float32): + input_tensor = Tensor(np.array( + [[[[0.1, 0.2, 0.3], [0.4, 0.5, 0.6], [0.7, 0.8, 0.9]]]]).astype(datatype)) + + # larger h and w + resize_nn = NetResizeBilinear((9, 9)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.13333334, 0.16666667, 0.2, 0.23333335, 0.26666668, 0.3, 0.3, 0.3], + [0.20000002, 0.23333335, 0.26666668, 0.3, 0.33333337, 0.3666667, 0.40000004, + 0.40000004, 0.40000004], + [0.3, 0.33333337, 0.36666667, 0.40000004, 0.43333337, 0.4666667, 0.5, 0.5, + 0.5], + [0.4, 0.43333334, 0.46666667, 0.5, 0.53333336, 0.5666667, 0.6, 0.6, 0.6], + [0.5, 0.53333336, 0.56666666, 0.6, 0.6333333, 0.66666675, 0.70000005, + 0.70000005, 0.70000005], + [0.6, 0.6333334, 0.6666667, 0.70000005, 0.73333335, 0.7666667, 0.8, 0.8, 0.8], + [0.7, 0.73333335, 0.76666665, 0.8, 0.8333333, 0.8666667, 0.9, 0.9, 0.9], + [0.7, 0.73333335, 0.76666665, 0.8, 0.8333333, 0.8666667, 0.9, 0.9, 0.9], + [0.7, 0.73333335, 0.76666665, 0.8, 0.8333333, 0.8666667, 0.9, 0.9, + 0.9]]]]).astype(np.float32)) + error = np.ones(shape=[9, 9]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h and w + resize_nn = NetResizeBilinear((1, 1)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1]]]]).astype(np.float32)) + error = np.ones(shape=[1, 1]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h, larger w + resize_nn = NetResizeBilinear((1, 6)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.1, 0.15, 0.2, 0.25, 0.3, 0.3]]]]).astype(np.float32)) + error = np.ones(shape=[1, 6]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # larger h, smaller w + resize_nn = NetResizeBilinear((6, 1)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.1], [0.25], [0.4], [0.55], [0.7], [0.7]]]]).astype(np.float32)) + error = np.ones(shape=[6, 1]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h, same w + resize_nn = NetResizeBilinear((1, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.1, 0.2, 0.3]]]]).astype(np.float32)) + error = np.ones(shape=[1, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # larger h, same w + resize_nn = NetResizeBilinear((6, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.2, 0.3], + [0.25, 0.35000002, 0.45000002], + [0.4, 0.5, 0.6], + [0.55, 0.65, 0.75], + [0.7, 0.8, 0.9], + [0.7, 0.8, 0.9]]]]).astype(np.float32)) + error = np.ones(shape=[6, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same h, smaller w + resize_nn = NetResizeBilinear((3, 1)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.1], [0.4], [0.7]]]]).astype(np.float32)) + error = np.ones(shape=[3, 1]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same h, larger w + resize_nn = NetResizeBilinear((3, 6)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.15, 0.2, 0.25, 0.3, 0.3], + [0.4, 0.45, 0.5, 0.55, 0.6, 0.6], + [0.7, 0.75, 0.8, 0.85, 0.9, 0.9]]]]).astype(np.float32)) + error = np.ones(shape=[3, 6]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same w, same h (identity) + resize_nn = NetResizeBilinear((3, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array( + [[[[0.1, 0.2, 0.3], [0.4, 0.5, 0.6], [0.7, 0.8, 0.9]]]]).astype(np.float32)) + error = np.ones(shape=[3, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + +@pytest.mark.level0 +@pytest.mark.platform_x86_gpu_training +@pytest.mark.env_onecard +def test_resize_nn_grayscale_not_integer_ratio_half(datatype=np.float16): + input_tensor = Tensor(np.array([[[[0.1, 0.2, 0.3, 0.4], + [0.5, 0.6, 0.7, 0.8], + [0.9, 0.0, 0.1, 0.2]]]]).astype(datatype)) + + # larger h and w + resize_nn = NetResizeBilinear((7, 7)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.15710449, 0.21425085, 0.2714495, 0.3285784, + 0.38563755, 0.39990234], + [0.27141464, 0.3285734, 0.3857422, 0.44294086, 0.5000399, + 0.55703926, 0.57128906], + [0.44285366, 0.5000423, 0.5572336, 0.6144322, 0.67150134, + 0.7284409, 0.7426758], + [0.6142578, 0.50819117, 0.44293588, 0.5001146, 0.5571937, + 0.6141731, 0.62841797], + [0.78564453, 0.4346799, 0.18574369, 0.2428925, 0.3000015, + 0.3570706, 0.3713379], + [0.89990234, 0.3856724, 0.01428223, 0.07141115, 0.12854005, + 0.18566895, 0.19995117], + [0.89990234, 0.3856724, 0.01428223, 0.07141115, 0.12854005, + 0.18566895, 0.19995117]]]]).astype(np.float32)) + error = np.ones(shape=[7, 7]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h and w + resize_nn = NetResizeBilinear((2, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.09997559, 0.23331706, 0.36661786], + [0.6999512, 0.33339438, 0.46661377]]]]).astype(np.float32)) + error = np.ones(shape=[2, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h, larger w + resize_nn = NetResizeBilinear((2, 7)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.15710449, 0.21425085, 0.2714495, 0.3285784, + 0.38563755, 0.39990234], + [0.6999512, 0.47143552, 0.3143398, 0.37150356, 0.4285976, + 0.48562187, 0.49987793]]]]).astype(np.float32)) + error = np.ones(shape=[2, 7]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # larger h, smaller w + resize_nn = NetResizeBilinear((5, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.23331706, 0.36661786], + [0.33999026, 0.47340494, 0.6066081], + [0.5799805, 0.51343584, 0.64660645], + [0.8199219, 0.15335283, 0.28662106], + [0.89990234, 0.0333252, 0.16662598]]]]).astype(np.float32)) + error = np.ones(shape=[5, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h, same w + resize_nn = NetResizeBilinear((2, 4)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.19995117, 0.30004883, 0.39990234], + [0.6999512, 0.30004883, 0.40008545, 0.49987793]]]]).astype(np.float32)) + error = np.ones(shape=[2, 4]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # larger h, same w + resize_nn = NetResizeBilinear((8, 4)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.19995117, 0.30004883, 0.39990234], + [0.24998474, 0.3500061, 0.45010376, 0.5498657], + [0.3999939, 0.50006104, 0.6001587, 0.6998291], + [0.5499878, 0.52508545, 0.62516785, 0.724823], + [0.6999512, 0.30004883, 0.40008545, 0.49987793], + [0.84991455, 0.07501221, 0.17500305, 0.27493286], + [0.89990234, 0., 0.09997559, 0.19995117], + [0.89990234, 0., 0.09997559, 0.19995117]]]]).astype(np.float32)) + error = np.ones(shape=[8, 4]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same h, smaller w + resize_nn = NetResizeBilinear((3, 2)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.30004883], + [0.5, 0.7001953], + [0.89990234, 0.09997559]]]]).astype(np.float32)) + error = np.ones(shape=[3, 2]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same h, larger w + resize_nn = NetResizeBilinear((3, 6)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.16662598, 0.23331706, 0.30004883, 0.36661786, + 0.39990234], + [0.5, 0.56673175, 0.63346356, 0.7001953, 0.76660156, + 0.7998047], + [0.89990234, 0.2999674, 0.0333252, 0.09997559, 0.16662598, + 0.19995117]]]]).astype(np.float32)) + error = np.ones(shape=[3, 6]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same w, same h (identity) + resize_nn = NetResizeBilinear((3, 4)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.19995117, 0.30004883, 0.39990234], + [0.5, 0.60009766, 0.7001953, 0.7998047], + [0.89990234, 0., 0.09997559, 0.19995117]]]]).astype(np.float32)) + error = np.ones(shape=[3, 4]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + +@pytest.mark.level0 +@pytest.mark.platform_x86_gpu_training +@pytest.mark.env_onecard +def test_resize_nn_grayscale_not_integer_ratio_float(datatype=np.float32): + input_tensor = Tensor(np.array([[[[0.1, 0.2, 0.3, 0.4], + [0.5, 0.6, 0.7, 0.8], + [0.9, 0.0, 0.1, 0.2]]]]).astype(datatype)) + + # larger h and w + resize_nn = NetResizeBilinear((7, 7)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.15714286, 0.21428573, 0.27142859, 0.32857144, 0.3857143, 0.4], + [0.27142859, 0.32857144, 0.38571432, 0.44285715, 0.5, 0.55714285, 0.5714286], + [0.44285715, 0.5, 0.5571429, 0.6142857, 0.67142856, 0.7285714, 0.74285716], + [0.6142857, 0.5081633, 0.4428572, 0.5, 0.55714285, 0.6142857, 0.62857145], + [0.78571427, 0.43469384, 0.1857143, 0.24285716, 0.3, 0.35714287, 0.37142855], + [0.9, 0.38571423, 0.01428572, 0.07142859, 0.12857144, 0.1857143, 0.2], + [0.9, 0.38571423, 0.01428572, 0.07142859, 0.12857144, 0.1857143, + 0.2]]]]).astype(np.float32)) + error = np.ones(shape=[7, 7]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h and w + resize_nn = NetResizeBilinear((2, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor( + np.array([[[[0.1, 0.23333335, 0.36666667], + [0.7, 0.33333334, 0.46666667]]]]).astype(np.float32)) + error = np.ones(shape=[2, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h, larger w + resize_nn = NetResizeBilinear((2, 7)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.15714286, 0.21428573, 0.27142859, 0.32857144, + 0.3857143, 0.4], + [0.7, 0.47142854, 0.31428576, 0.37142858, 0.42857143, + 0.4857143, 0.5]]]]).astype(np.float32)) + error = np.ones(shape=[2, 7]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # larger h, smaller w + resize_nn = NetResizeBilinear((5, 3)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.23333335, 0.36666667], + [0.34, 0.47333336, 0.6066667], + [0.58000004, 0.5133333, 0.64666665], + [0.82000005, 0.1533333, 0.28666663], + [0.9, 0.03333334, 0.16666669]]]]).astype(np.float32)) + error = np.ones(shape=[5, 3]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # smaller h, same w + resize_nn = NetResizeBilinear((2, 4)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.2, 0.3, 0.4], + [0.7, 0.3, 0.4, 0.5]]]]).astype(np.float32)) + error = np.ones(shape=[2, 4]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # larger h, same w + resize_nn = NetResizeBilinear((8, 4)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.2, 0.3, 0.4], + [0.25, 0.35000002, 0.45, 0.55], + [0.4, 0.5, 0.6, 0.70000005], + [0.55, 0.52500004, 0.625, 0.725], + [0.7, 0.3, 0.4, 0.5], + [0.84999996, 0.07499999, + 0.17500001, 0.27499998], + [0.9, 0., 0.1, 0.2], + [0.9, 0., 0.1, 0.2]]]]).astype(np.float32)) + error = np.ones(shape=[8, 4]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same h, smaller w + resize_nn = NetResizeBilinear((3, 2)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.3], + [0.5, 0.7], + [0.9, 0.1]]]]).astype(np.float32)) + error = np.ones(shape=[3, 2]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same h, larger w + resize_nn = NetResizeBilinear((3, 6)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.16666667, 0.23333335, 0.3, 0.36666667, 0.4], + [0.5, 0.56666666, 0.6333333, 0.7, 0.76666665, 0.8], + [0.9, 0.29999995, 0.03333334, 0.1, 0.16666669, 0.2]]]]).astype(np.float32)) + error = np.ones(shape=[3, 6]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + # same w, same h (identity) + resize_nn = NetResizeBilinear((3, 4)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.1, 0.2, 0.3, 0.4], + [0.5, 0.6, 0.7, 0.8], + [0.9, 0., 0.1, 0.2]]]]).astype(np.float32)) + error = np.ones(shape=[3, 4]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + +@pytest.mark.level0 +@pytest.mark.platform_x86_gpu_training +@pytest.mark.env_onecard +def test_resize_nn_grayscale_multiple_images_half(datatype=np.float16): + input_tensor = Tensor(np.array([[[[0.1, 0.2, 0.3], [0.4, 0.5, 0.6], [0.7, 0.8, 0.9]]], + [[[0.4, 0.5, 0.6], [0.7, 0.8, 0.9], [0.1, 0.2, 0.3]]], + [[[0.7, 0.8, 0.9], [0.1, 0.2, 0.3], [0.4, 0.5, 0.6]]]]).astype(datatype)) + + resize_nn = NetResizeBilinear((2, 6)) + output = resize_nn(input_tensor) + expected_output = Tensor(np.array([[[[0.09997559, 0.14996338, 0.19995117, 0.25, 0.30004883, 0.30004883], + [0.5500488, 0.5999756, 0.64990234, 0.6999512, 0.75, 0.75]]], + [[[0.39990234, 0.44995117, 0.5, 0.5500488, 0.60009766, 0.60009766], + [0.40008545, 0.4499817, 0.49987793, 0.54992676, 0.5999756, 0.5999756]]], + [[[0.7001953, 0.75, 0.7998047, 0.8498535, 0.89990234, 0.89990234], + [0.24993896, 0.29995728, 0.3499756, 0.4000244, 0.45007324, + 0.45007324]]]]).astype(np.float32)) + + error = np.ones(shape=[3, 3, 2, 6]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + +@pytest.mark.level0 +@pytest.mark.platform_x86_gpu_training +@pytest.mark.env_onecard +def test_resize_nn_grayscale_multiple_images_float(datatype=np.float32): + input_tensor = Tensor(np.array([[[[0.1, 0.2, 0.3], [0.4, 0.5, 0.6], [0.7, 0.8, 0.9]]], + [[[0.4, 0.5, 0.6], [0.7, 0.8, 0.9], [0.1, 0.2, 0.3]]], + [[[0.7, 0.8, 0.9], [0.1, 0.2, 0.3], [0.4, 0.5, 0.6]]]]).astype(datatype)) + + resize_nn = NetResizeBilinear((2, 6)) + output = resize_nn(input_tensor) + + expected_output = Tensor(np.array([[[[0.1, 0.15, 0.2, 0.25, 0.3, 0.3], + [0.55, 0.6, 0.65, 0.70000005, 0.75, 0.75]]], + [[[0.4, 0.45, 0.5, 0.55, 0.6, 0.6], + [0.4, 0.45, 0.5, 0.55, 0.6, 0.6]]], + [[[0.7, 0.75, 0.8, 0.85, 0.9, 0.9], + [0.25, 0.3, 0.35000002, 0.4, 0.45000002, 0.45000002]]]]).astype(np.float32)) + + error = np.ones(shape=[3, 3, 2, 6]) * 1.0e-6 + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + + +@pytest.mark.level0 +@pytest.mark.platform_x86_gpu_training +@pytest.mark.env_onecard +def test_resize_nn_grayscale_align_corners_half(datatype=np.float16): + input_tensor = Tensor( + np.array([[[[0.1, 0.2, 0.3, 0.4], [0.5, 0.6, 0.7, 0.8]]]]).astype(datatype)) + + resize_nn_corners_aligned = NetResizeBilinear( + size=(3, 7), align_corner=True) + output_corners_aligned = resize_nn_corners_aligned(input_tensor) + + resize_nn = NetResizeBilinear((3, 7)) + output = resize_nn(input_tensor) + + expected_output_align = Tensor(np.array([[[[0.09997559, 0.14996338, 0.19995117, 0.25, 0.30004883, + 0.3499756, 0.39990234], + [0.2999878, 0.3500061, 0.4000244, 0.45007324, 0.5001221, + 0.5499878, 0.5998535], + [0.5, 0.5500488, 0.60009766, 0.6501465, 0.7001953, + 0.75, 0.7998047]]]]).astype(np.float32)) + expected_output = Tensor(np.array([[[[0.09997559, 0.15710449, 0.21425085, 0.2714495, 0.3285784, + 0.38563755, 0.39990234], + [0.36665854, 0.42383394, 0.4810152, 0.53821385, 0.59529626, + 0.6522624, 0.6665039], + [0.5, 0.55719864, 0.61439735, 0.671596, 0.72865516, + 0.7855748, 0.7998047]]]]).astype(np.float32)) + + error = np.ones(shape=[3, 7]) * 1.0e-6 + diff_align = output_corners_aligned.asnumpy() - expected_output_align.asnumpy() + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + assert np.all(abs(diff_align) < error) + + +@pytest.mark.level0 +@pytest.mark.platform_x86_gpu_training +@pytest.mark.env_onecard +def test_resize_nn_grayscale_align_corners_float(datatype=np.float32): + input_tensor = Tensor( + np.array([[[[0.1, 0.2, 0.3, 0.4], [0.5, 0.6, 0.7, 0.8]]]]).astype(datatype)) + + resize_nn_corners_aligned = NetResizeBilinear( + size=(3, 7), align_corner=True) + output_corners_aligned = resize_nn_corners_aligned(input_tensor) + + resize_nn = NetResizeBilinear((3, 7)) + output = resize_nn(input_tensor) + + expected_output_align = Tensor(np.array([[[[0.1, 0.15, 0.2, 0.25, 0.3, + 0.35000002, 0.4], + [0.3, 0.35000002, 0.40000004, 0.45, 0.5, + 0.55, 0.6], + [0.5, 0.55, 0.6, 0.65, 0.7, + 0.75, 0.8]]]]).astype(datatype)) + expected_output = Tensor(np.array([[[[0.1, 0.15714286, 0.21428573, 0.27142859, 0.32857144, + 0.3857143, 0.4], + [0.36666667, 0.42380953, 0.48095244, 0.53809524, 0.5952381, + 0.65238094, 0.6666667], + [0.5, 0.55714285, 0.61428577, 0.67142856, 0.7285714, + 0.78571427, 0.8]]]]).astype(datatype)) + + error = np.ones(shape=[3, 7]) * 1.0e-6 + diff_align = output_corners_aligned.asnumpy() - expected_output_align.asnumpy() + diff = output.asnumpy() - expected_output.asnumpy() + assert np.all(abs(diff) < error) + assert np.all(abs(diff_align) < error)