experiments/STT-tensorflow
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Resize task modified to be Metal compatible.

Browse Source 
Added resize_test_util with unified tests.
Added Metal resize unit test.
Added new defines.

PiperOrigin-RevId: 353172824
Change-Id: I0a3b575a9e9738a6486a10d2ece0b70926efb93d
This commit is contained in:
Raman Sarokin 2021-01-21 22:08:44 -08:00 committed by TensorFlower Gardener
parent 43cbed5c96
commit ab9eb3d104
10 changed files with 396 additions and 242 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
tensorflow/lite/delegates/gpu/cl/cl_operation.cc
View File

@ -60,14 +60,17 @@ std::string GetCommonOpenCLDefines(CalculationsPrecision precision) {
result += "#define GROUP_ID_2 get_group_id(2)\n";
result += "#define LOCAL_MEM_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n";
result += "#define MAIN_FUNCTION __kernel void main_function\n";
result += "#define INIT_FLOAT(value) (float)(value)\n";
result += "#define INIT_FLOAT2(value) (float2)(value)\n";
result += "#define INIT_FLOAT2v2(v0, v1) (float2)(v0, v1)\n";
result += "#define INIT_FLOAT3(value) (float3)(value)\n";
result += "#define INIT_FLOAT3v3(v0, v1, v2) (float3)(v0, v1, v2)\n";
result += "#define INIT_FLOAT4(value) (float4)(value)\n";
result += "#define INIT_FLOAT4v4(v0, v1, v2, v3) (float4)(v0, v1, v2, v3)\n";
result += "#define INIT_INT(value) (int)(value)\n";
result += "#define INIT_INT2v2(v0, v1) (int2)(v0, v1)\n";
result += "#define INIT_INT4v4(v0, v1, v2, v3) (int4)(v0, v1, v2, v3)\n";
result += "#define CONVERT_TO_INT4(value) convert_int4(value)\n";
switch (precision) {
case CalculationsPrecision::F32:
result += "#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n";

2
tensorflow/lite/delegates/gpu/cl/kernels/BUILD
View File

@ -596,7 +596,7 @@ cc_test(
":cl_test",
"//tensorflow/lite/delegates/gpu/common:operations",
"//tensorflow/lite/delegates/gpu/common:status",
"//tensorflow/lite/delegates/gpu/common/tasks:resize",
"//tensorflow/lite/delegates/gpu/common/tasks:resize_test_util",
"@com_google_googletest//:gtest_main",
],
)

226
tensorflow/lite/delegates/gpu/cl/kernels/resize_test.cc
View File

@ -13,8 +13,6 @@ See the License for the specific language governing permissions and
limitations under the License.
==============================================================================*/
#include "tensorflow/lite/delegates/gpu/common/tasks/resize.h"
#include <vector>
#include <gmock/gmock.h>
@ -22,9 +20,7 @@ limitations under the License.
#include "tensorflow/lite/delegates/gpu/cl/kernels/cl_test.h"
#include "tensorflow/lite/delegates/gpu/common/operations.h"
#include "tensorflow/lite/delegates/gpu/common/status.h"
using ::testing::FloatNear;
using ::testing::Pointwise;
#include "tensorflow/lite/delegates/gpu/common/tasks/resize_test_util.h"
namespace tflite {
namespace gpu {
@ -32,230 +28,38 @@ namespace cl {
namespace {
TEST_F(OpenCLOperationTest, ResizeBilinearAligned) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 3, 1);
src_tensor.data = {0.0f, 1.0f, 2.0f, 3.0f, 4.0f, 5.0f};
Resize2DAttributes attr;
attr.type = SamplingType::BILINEAR;
attr.new_shape = HW(4, 4);
attr.align_corners = true;
for (auto storage : env_.GetSupportedStorages()) {
for (auto precision : env_.GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
ASSERT_OK(
ExecuteGPUOperation(src_tensor, creation_context_,
absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 4, 4, 1), &dst_tensor));
EXPECT_THAT(dst_tensor.data,
Pointwise(FloatNear(eps),
{0.0f, 0.666667f, 1.33333f, 2.0f, 1.0f, 1.66667f,
2.33333f, 3.0f, 2.0f, 2.66667f, 3.33333f, 4.0f,
3.0f, 3.66667f, 4.33333f, 5.0f}));
}
}
auto status = ResizeBilinearAlignedTest(&exec_env_);
ASSERT_TRUE(status.ok()) << status.error_message();
}
TEST_F(OpenCLOperationTest, ResizeBilinearNonAligned) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 3, 1);
src_tensor.data = {0.0f, 1.0f, 2.0f, 3.0f, 4.0f, 5.0f};
Resize2DAttributes attr;
attr.type = SamplingType::BILINEAR;
attr.new_shape = HW(4, 4);
attr.align_corners = false;
for (auto storage : env_.GetSupportedStorages()) {
for (auto precision : env_.GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
ASSERT_OK(
ExecuteGPUOperation(src_tensor, creation_context_,
absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 4, 4, 1), &dst_tensor));
EXPECT_THAT(
dst_tensor.data,
Pointwise(FloatNear(eps),
{0.0f, 0.75f, 1.5f, 2.0f, 1.5f, 2.25f, 3.0f, 3.5f, 3.0f,
3.75f, 4.5f, 5.0f, 3.0f, 3.75f, 4.5f, 5.0f}));
}
}
auto status = ResizeBilinearNonAlignedTest(&exec_env_);
ASSERT_TRUE(status.ok()) << status.error_message();
}
TEST_F(OpenCLOperationTest, ResizeBilinearWithoutHalfPixel) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 2, 1);
src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f};
Resize2DAttributes attr;
attr.type = SamplingType::BILINEAR;
attr.new_shape = HW(3, 3);
attr.align_corners = false;
attr.half_pixel_centers = false;
for (auto storage : env_.GetSupportedStorages()) {
for (auto precision : env_.GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
ASSERT_OK(
ExecuteGPUOperation(src_tensor, creation_context_,
absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 3, 3, 1), &dst_tensor));
EXPECT_THAT(
dst_tensor.data,
Pointwise(FloatNear(eps), {1.0f, 1.666666f, 2.0f, 2.333333f, 3.0f,
3.333333f, 3.0f, 3.666666f, 4.0f}));
}
}
auto status = ResizeBilinearWithoutHalfPixelTest(&exec_env_);
ASSERT_TRUE(status.ok()) << status.error_message();
}
TEST_F(OpenCLOperationTest, ResizeBilinearWithHalfPixel) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 2, 1);
src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f};
Resize2DAttributes attr;
attr.type = SamplingType::BILINEAR;
attr.new_shape = HW(3, 3);
attr.align_corners = false;
attr.half_pixel_centers = true;
for (auto storage : env_.GetSupportedStorages()) {
for (auto precision : env_.GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
ASSERT_OK(
ExecuteGPUOperation(src_tensor, creation_context_,
absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 3, 3, 1), &dst_tensor));
EXPECT_THAT(dst_tensor.data,
Pointwise(FloatNear(eps), {1.0f, 1.5f, 2.0f, 2.0f, 2.5f, 3.0f,
3.0f, 3.5f, 4.0f}));
}
}
auto status = ResizeBilinearWithHalfPixelTest(&exec_env_);
ASSERT_TRUE(status.ok()) << status.error_message();
}
TEST_F(OpenCLOperationTest, ResizeNearest) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 1, 2, 1);
src_tensor.data = {1.0f, 2.0f};
Resize2DAttributes attr;
attr.align_corners = false;
attr.half_pixel_centers = false;
attr.new_shape = HW(2, 4);
attr.type = SamplingType::NEAREST;
for (auto storage : env_.GetSupportedStorages()) {
for (auto precision : env_.GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
ASSERT_OK(
ExecuteGPUOperation(src_tensor, creation_context_,
absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 2, 4, 1), &dst_tensor));
EXPECT_THAT(dst_tensor.data,
Pointwise(FloatNear(eps),
{1.0f, 1.0f, 2.0f, 2.0f, 1.0f, 1.0f, 2.0f, 2.0f}));
}
}
auto status = ResizeNearestTest(&exec_env_);
ASSERT_TRUE(status.ok()) << status.error_message();
}
TEST_F(OpenCLOperationTest, ResizeNearestAlignCorners) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 2, 1);
src_tensor.data = {3.0f, 6.0f, 9.0f, 12.0f};
Resize2DAttributes attr;
attr.align_corners = true;
attr.half_pixel_centers = false;
attr.new_shape = HW(3, 3);
attr.type = SamplingType::NEAREST;
for (auto storage : env_.GetSupportedStorages()) {
for (auto precision : env_.GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
ASSERT_OK(
ExecuteGPUOperation(src_tensor, creation_context_,
absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 3, 3, 1), &dst_tensor));
EXPECT_THAT(dst_tensor.data,
Pointwise(FloatNear(eps), {3.0f, 6.0f, 6.0f, 9.0f, 12.0f,
12.0f, 9.0f, 12.0f, 12.0f}));
}
}
auto status = ResizeNearestAlignCornersTest(&exec_env_);
ASSERT_TRUE(status.ok()) << status.error_message();
}
TEST_F(OpenCLOperationTest, ResizeNearestHalfPixelCenters) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 2, 1);
src_tensor.data = {3.0f, 6.0f, 9.0f, 12.0f};
Resize2DAttributes attr;
attr.align_corners = false;
attr.half_pixel_centers = true;
attr.new_shape = HW(3, 3);
attr.type = SamplingType::NEAREST;
for (auto storage : env_.GetSupportedStorages()) {
for (auto precision : env_.GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
ASSERT_OK(
ExecuteGPUOperation(src_tensor, creation_context_,
absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 3, 3, 1), &dst_tensor));
EXPECT_THAT(dst_tensor.data,
Pointwise(FloatNear(eps), {3.0f, 6.0f, 6.0f, 9.0f, 12.0f,
12.0f, 9.0f, 12.0f, 12.0f}));
}
}
auto status = ResizeNearestHalfPixelCentersTest(&exec_env_);
ASSERT_TRUE(status.ok()) << status.error_message();
}
} // namespace

13
tensorflow/lite/delegates/gpu/common/tasks/BUILD
View File

@ -613,6 +613,19 @@ cc_library(
],
)
cc_library(
name = "resize_test_util",
testonly = 1,
srcs = ["resize_test_util.cc"],
hdrs = ["resize_test_util.h"],
deps = [
":resize",
"//tensorflow/lite/delegates/gpu/common:operations",
"//tensorflow/lite/delegates/gpu/common:status",
"//tensorflow/lite/delegates/gpu/common/task:testing_util",
],
)
cc_library(
name = "softmax",
srcs = ["softmax.cc"],

55
tensorflow/lite/delegates/gpu/common/tasks/resize.cc
View File

@ -55,18 +55,17 @@ std::string Resize::GetResizeCode(const OperationDef& op_def,
args_.AddFloat("scale_factor_y");
std::string c;
c += "__kernel void main_function(\n";
c += "$0) {\n";
c += " int Y = get_global_id(1);\n";
c += " int Z = get_global_id(2);\n";
c += "MAIN_FUNCTION($0) {\n";
c += " int Y = GLOBAL_ID_1;\n";
c += " int Z = GLOBAL_ID_2;\n";
if (op_def.IsBatchSupported()) {
c += " int linear_id = get_global_id(0);\n";
c += " int linear_id = GLOBAL_ID_0;\n";
c += " int X = linear_id / args.dst_tensor.Batch();\n";
c += " int B = linear_id % args.dst_tensor.Batch();\n";
c += " if (linear_id >= args.dst_tensor.Width() || Y >= "
"args.dst_tensor.Height() || Z >= args.dst_tensor.Slices()) return;\n";
} else {
c += " int X = get_global_id(0);\n";
c += " int X = GLOBAL_ID_0;\n";
c += " if (X >= args.dst_tensor.Width() || Y >= args.dst_tensor.Height() "
"|| Z >= args.dst_tensor.Slices()) return;\n";
}
@ -85,8 +84,8 @@ std::string Resize::GetResizeCode(const OperationDef& op_def,
fyc += " + 0.5f";
}
c += " int2 coord;\n";
c += " coord.x = (int)(" + fxc + ");\n";
c += " coord.y = (int)(" + fyc + ");\n";
c += " coord.x = INIT_INT(" + fxc + ");\n";
c += " coord.y = INIT_INT(" + fyc + ");\n";
c += " coord.x = max(0, coord.x);\n";
c += " coord.y = max(0, coord.y);\n";
c += " coord.x = min(coord.x, args.border_x);\n";
@ -98,18 +97,21 @@ std::string Resize::GetResizeCode(const OperationDef& op_def,
c += " FLT4 r0 = args.src_tensor.Read(coord.x, coord.y, Z);\n";
} else {
if (attr.half_pixel_centers) {
c += " float2 f_coords = ((float2)(X, Y) + 0.5f) * "
"(float2)(args.scale_factor_x, args.scale_factor_y) - "
c += " float2 f_coords = (INIT_FLOAT2v2(X, Y) + 0.5f) * "
"INIT_FLOAT2v2(args.scale_factor_x, args.scale_factor_y) - "
"0.5f;\n";
} else {
c += " float2 f_coords = (float2)(X, Y) * (float2)(args.scale_factor_x, "
c += " float2 f_coords = INIT_FLOAT2v2(X, Y) * "
"INIT_FLOAT2v2(args.scale_factor_x, "
"args.scale_factor_y);\n";
}
c += " float2 f_coords_floor = floor(f_coords);\n";
c += " int2 coords_floor = (int2)(f_coords_floor.x, f_coords_floor.y);\n";
c += " int2 coords_floor = INIT_INT2v2(f_coords_floor.x, "
"f_coords_floor.y);\n";
c += " int4 st;\n";
c += " st.xy = max(coords_floor, (int2)(0, 0));\n";
c += " st.zw = min(coords_floor + (int2)(1, 1), (int2)(args.border_x, "
c += " st.xy = max(coords_floor, INIT_INT2v2(0, 0));\n";
c += " st.zw = min(coords_floor + INIT_INT2v2(1, 1), "
"INIT_INT2v2(args.border_x, "
"args.border_y));\n";
c += " float2 t = f_coords - f_coords_floor;\n";
if (op_def.IsBatchSupported()) {
@ -190,20 +192,19 @@ std::string Resize3D::GetResize3DCode(const OperationDef& op_def,
args_.AddFloat("scale_factor_z");
std::string c;
c += "__kernel void main_function(\n";
c += "$0) {\n";
c += " int Y = get_global_id(1);\n";
c += " int linear_id_z = get_global_id(2);\n";
c += "MAIN_FUNCTION($0) {\n";
c += " int Y = GLOBAL_ID_1;\n";
c += " int linear_id_z = GLOBAL_ID_2;\n";
c += " int S = linear_id_z % args.dst_tensor.Slices();\n";
c += " int Z = linear_id_z / args.dst_tensor.Slices();\n";
if (op_def.IsBatchSupported()) {
c += " int linear_id = get_global_id(0);\n";
c += " int linear_id = GLOBAL_ID_0;\n";
c += " int X = linear_id / args.dst_tensor.Batch();\n";
c += " int B = linear_id % args.dst_tensor.Batch();\n";
c += " if (linear_id >= args.dst_tensor.Width() || Y >= "
"args.dst_tensor.Height() || Z >= args.dst_tensor.Depth()) return;\n";
} else {
c += " int X = get_global_id(0);\n";
c += " int X = GLOBAL_ID_0;\n";
c += " if (X >= args.dst_tensor.Width() || Y >= args.dst_tensor.Height() "
"|| Z >= args.dst_tensor.Depth()) return;\n";
}
@ -226,9 +227,9 @@ std::string Resize3D::GetResize3DCode(const OperationDef& op_def,
fzc += " + 0.5f";
}
c += " int4 coord;\n";
c += " coord.x = (int)(" + fxc + ");\n";
c += " coord.y = (int)(" + fyc + ");\n";
c += " coord.z = (int)(" + fzc + ");\n";
c += " coord.x = INIT_INT(" + fxc + ");\n";
c += " coord.y = INIT_INT(" + fyc + ");\n";
c += " coord.z = INIT_INT(" + fzc + ");\n";
c += " coord.x = max(0, coord.x);\n";
c += " coord.y = max(0, coord.y);\n";
c += " coord.z = max(0, coord.z);\n";
@ -242,10 +243,10 @@ std::string Resize3D::GetResize3DCode(const OperationDef& op_def,
c += " FLT4 r0 = args.src_tensor.Read(coord.x, coord.y, coord.z, S);\n";
} else {
c += " float4 f_coords;\n";
c += " f_coords.x = (float)(X) * args.scale_factor_x;\n";
c += " f_coords.y = (float)(Y) * args.scale_factor_y;\n";
c += " f_coords.z = (float)(Z) * args.scale_factor_z;\n";
c += " int4 start = (int4)(f_coords.x, f_coords.y, f_coords.z, 0);\n";
c += " f_coords.x = INIT_FLOAT(X) * args.scale_factor_x;\n";
c += " f_coords.y = INIT_FLOAT(Y) * args.scale_factor_y;\n";
c += " f_coords.z = INIT_FLOAT(Z) * args.scale_factor_z;\n";
c += " int4 start = INIT_INT4v4(f_coords.x, f_coords.y, f_coords.z, 0);\n";
c += " int4 end;\n";
c += " end.x = min(start.x + 1, args.border_x);\n";
c += " end.y = min(start.y + 1, args.border_y);\n";

253
tensorflow/lite/delegates/gpu/common/tasks/resize_test_util.cc Normal file
View File

@ -0,0 +1,253 @@
/* Copyright 2021 The TensorFlow Authors. All Rights Reserved.
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 "tensorflow/lite/delegates/gpu/common/tasks/resize_test_util.h"
#include <vector>
#include "tensorflow/lite/delegates/gpu/common/operations.h"
#include "tensorflow/lite/delegates/gpu/common/status.h"
#include "tensorflow/lite/delegates/gpu/common/task/testing_util.h"
#include "tensorflow/lite/delegates/gpu/common/tasks/resize.h"
namespace tflite {
namespace gpu {
absl::Status ResizeBilinearAlignedTest(TestExecutionEnvironment* env) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 3, 1);
src_tensor.data = {0.0f, 1.0f, 2.0f, 3.0f, 4.0f, 5.0f};
Resize2DAttributes attr;
attr.type = SamplingType::BILINEAR;
attr.new_shape = HW(4, 4);
attr.align_corners = true;
for (auto storage : env->GetSupportedStorages()) {
for (auto precision : env->GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
RETURN_IF_ERROR(env->ExecuteGPUOperation(
src_tensor, absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 4, 4, 1), &dst_tensor));
RETURN_IF_ERROR(PointWiseNear(
{0.0f, 0.666667f, 1.33333f, 2.0f, 1.0f, 1.66667f, 2.33333f, 3.0f,
2.0f, 2.66667f, 3.33333f, 4.0f, 3.0f, 3.66667f, 4.33333f, 5.0f},
dst_tensor.data, eps));
}
}
return absl::OkStatus();
}
absl::Status ResizeBilinearNonAlignedTest(TestExecutionEnvironment* env) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 3, 1);
src_tensor.data = {0.0f, 1.0f, 2.0f, 3.0f, 4.0f, 5.0f};
Resize2DAttributes attr;
attr.type = SamplingType::BILINEAR;
attr.new_shape = HW(4, 4);
attr.align_corners = false;
for (auto storage : env->GetSupportedStorages()) {
for (auto precision : env->GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
RETURN_IF_ERROR(env->ExecuteGPUOperation(
src_tensor, absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 4, 4, 1), &dst_tensor));
RETURN_IF_ERROR(
PointWiseNear({0.0f, 0.75f, 1.5f, 2.0f, 1.5f, 2.25f, 3.0f, 3.5f, 3.0f,
3.75f, 4.5f, 5.0f, 3.0f, 3.75f, 4.5f, 5.0f},
dst_tensor.data, eps));
}
}
return absl::OkStatus();
}
absl::Status ResizeBilinearWithoutHalfPixelTest(TestExecutionEnvironment* env) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 2, 1);
src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f};
Resize2DAttributes attr;
attr.type = SamplingType::BILINEAR;
attr.new_shape = HW(3, 3);
attr.align_corners = false;
attr.half_pixel_centers = false;
for (auto storage : env->GetSupportedStorages()) {
for (auto precision : env->GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
RETURN_IF_ERROR(env->ExecuteGPUOperation(
src_tensor, absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 3, 3, 1), &dst_tensor));
RETURN_IF_ERROR(PointWiseNear({1.0f, 1.666666f, 2.0f, 2.333333f, 3.0f,
3.333333f, 3.0f, 3.666666f, 4.0f},
dst_tensor.data, eps));
}
}
return absl::OkStatus();
}
absl::Status ResizeBilinearWithHalfPixelTest(TestExecutionEnvironment* env) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 2, 1);
src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f};
Resize2DAttributes attr;
attr.type = SamplingType::BILINEAR;
attr.new_shape = HW(3, 3);
attr.align_corners = false;
attr.half_pixel_centers = true;
for (auto storage : env->GetSupportedStorages()) {
for (auto precision : env->GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
RETURN_IF_ERROR(env->ExecuteGPUOperation(
src_tensor, absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 3, 3, 1), &dst_tensor));
RETURN_IF_ERROR(
PointWiseNear({1.0f, 1.5f, 2.0f, 2.0f, 2.5f, 3.0f, 3.0f, 3.5f, 4.0f},
dst_tensor.data, eps));
}
}
return absl::OkStatus();
}
absl::Status ResizeNearestTest(TestExecutionEnvironment* env) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 1, 2, 1);
src_tensor.data = {1.0f, 2.0f};
Resize2DAttributes attr;
attr.align_corners = false;
attr.half_pixel_centers = false;
attr.new_shape = HW(2, 4);
attr.type = SamplingType::NEAREST;
for (auto storage : env->GetSupportedStorages()) {
for (auto precision : env->GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
RETURN_IF_ERROR(env->ExecuteGPUOperation(
src_tensor, absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 2, 4, 1), &dst_tensor));
RETURN_IF_ERROR(
PointWiseNear({1.0f, 1.0f, 2.0f, 2.0f, 1.0f, 1.0f, 2.0f, 2.0f},
dst_tensor.data, eps));
}
}
return absl::OkStatus();
}
absl::Status ResizeNearestAlignCornersTest(TestExecutionEnvironment* env) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 2, 1);
src_tensor.data = {3.0f, 6.0f, 9.0f, 12.0f};
Resize2DAttributes attr;
attr.align_corners = true;
attr.half_pixel_centers = false;
attr.new_shape = HW(3, 3);
attr.type = SamplingType::NEAREST;
for (auto storage : env->GetSupportedStorages()) {
for (auto precision : env->GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
RETURN_IF_ERROR(env->ExecuteGPUOperation(
src_tensor, absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 3, 3, 1), &dst_tensor));
RETURN_IF_ERROR(PointWiseNear(
{3.0f, 6.0f, 6.0f, 9.0f, 12.0f, 12.0f, 9.0f, 12.0f, 12.0f},
dst_tensor.data, eps));
}
}
return absl::OkStatus();
}
absl::Status ResizeNearestHalfPixelCentersTest(TestExecutionEnvironment* env) {
TensorFloat32 src_tensor;
src_tensor.shape = BHWC(1, 2, 2, 1);
src_tensor.data = {3.0f, 6.0f, 9.0f, 12.0f};
Resize2DAttributes attr;
attr.align_corners = false;
attr.half_pixel_centers = true;
attr.new_shape = HW(3, 3);
attr.type = SamplingType::NEAREST;
for (auto storage : env->GetSupportedStorages()) {
for (auto precision : env->GetSupportedPrecisions()) {
const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 1e-2f;
OperationDef op_def;
op_def.precision = precision;
auto data_type = DeduceDataTypeFromPrecision(precision);
op_def.src_tensors.push_back({data_type, storage, Layout::HWC});
op_def.dst_tensors.push_back({data_type, storage, Layout::HWC});
TensorFloat32 dst_tensor;
Resize operation = CreateResize(op_def, attr);
RETURN_IF_ERROR(env->ExecuteGPUOperation(
src_tensor, absl::make_unique<Resize>(std::move(operation)),
BHWC(1, 3, 3, 1), &dst_tensor));
RETURN_IF_ERROR(PointWiseNear(
{3.0f, 6.0f, 6.0f, 9.0f, 12.0f, 12.0f, 9.0f, 12.0f, 12.0f},
dst_tensor.data, eps));
}
}
return absl::OkStatus();
}
} // namespace gpu
} // namespace tflite

36
tensorflow/lite/delegates/gpu/common/tasks/resize_test_util.h Normal file
View File

@ -0,0 +1,36 @@
/* Copyright 2021 The TensorFlow Authors. All Rights Reserved.
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 TENSORFLOW_LITE_DELEGATES_GPU_COMMON_TASKS_RESIZE_TEST_UTIL_H_
#define TENSORFLOW_LITE_DELEGATES_GPU_COMMON_TASKS_RESIZE_TEST_UTIL_H_
#include "tensorflow/lite/delegates/gpu/common/status.h"
#include "tensorflow/lite/delegates/gpu/common/task/testing_util.h"
namespace tflite {
namespace gpu {
absl::Status ResizeBilinearAlignedTest(TestExecutionEnvironment* env);
absl::Status ResizeBilinearNonAlignedTest(TestExecutionEnvironment* env);
absl::Status ResizeBilinearWithoutHalfPixelTest(TestExecutionEnvironment* env);
absl::Status ResizeBilinearWithHalfPixelTest(TestExecutionEnvironment* env);
absl::Status ResizeNearestTest(TestExecutionEnvironment* env);
absl::Status ResizeNearestAlignCornersTest(TestExecutionEnvironment* env);
absl::Status ResizeNearestHalfPixelCentersTest(TestExecutionEnvironment* env);
} // namespace gpu
} // namespace tflite
#endif // TENSORFLOW_LITE_DELEGATES_GPU_COMMON_TASKS_RESIZE_TEST_UTIL_H_

3
tensorflow/lite/delegates/gpu/metal/compute_task.cc
View File

@ -200,14 +200,17 @@ absl::Status ComputeTask::CompileProgram(MetalDevice* device,
[NSString stringWithFormat:@"%@4(value)", storageType],
@"\"INIT_FLT4v4(v0, v1, v2, v3)\"" :
[NSString stringWithFormat:@"\"%@4(v0, v1, v2, v3)\"", storageType],
@"INIT_FLOAT(value)" : @"float(value)",
@"INIT_FLOAT2(value)" : @"float2(value)",
@"\"INIT_FLOAT2v2(v0, v1)\"" : @"\"float2(v0, v1)\"",
@"INIT_FLOAT3(value)" : @"float3(value)",
@"\"INIT_FLOAT3v3(v0, v1, v2)\"" : @"\"float3(v0, v1, v2)\"",
@"INIT_FLOAT4(value)" : @"float4(value)",
@"\"INIT_FLOAT4v4(v0, v1, v2, v3)\"" : @"\"float4(v0, v1, v2, v3)\"",
@"INIT_INT(value)" : @"int(value)",
@"\"INIT_INT2v2(v0, v1)\"" : @"\"int2(v0, v1)\"",
@"\"INIT_INT4v4(v0, v1, v2, v3)\"" : @"\"int4(v0, v1, v2, v3)\"",
@"CONVERT_TO_INT4(value)" : @"int4(value)",
};
NSString* code =

2
tensorflow/lite/delegates/gpu/metal/kernels/BUILD
View File

@ -511,6 +511,7 @@ objc_library(
deps = [
":resize",
":test_util",
"//tensorflow/lite/delegates/gpu/common/tasks:resize_test_util",
],
)
@ -829,6 +830,7 @@ objc_library(
"//tensorflow/lite/delegates/gpu/common/tasks:reduce_test_util",
"//tensorflow/lite/delegates/gpu/common/tasks:relu_test_util",
"//tensorflow/lite/delegates/gpu/common/tasks:reshape_test_util",
"//tensorflow/lite/delegates/gpu/common/tasks:resize_test_util",
"//tensorflow/lite/delegates/gpu/common/tasks:softmax_test_util",
"//tensorflow/lite/delegates/gpu/common/tasks:space_to_depth_test_util",
"//tensorflow/lite/delegates/gpu/common/tasks:strided_slice_test_util",

45
tensorflow/lite/delegates/gpu/metal/kernels/resize_test.mm
View File

@ -23,6 +23,7 @@ limitations under the License.
#include "tensorflow/lite/delegates/gpu/common/operations.h"
#include "tensorflow/lite/delegates/gpu/common/shape.h"
#include "tensorflow/lite/delegates/gpu/common/status.h"
#include "tensorflow/lite/delegates/gpu/common/tasks/resize_test_util.h"
#include "tensorflow/lite/delegates/gpu/common/tensor.h"
#include "tensorflow/lite/delegates/gpu/common/util.h"
#include "tensorflow/lite/delegates/gpu/metal/compute_task_descriptor.h"
@ -41,7 +42,10 @@ using ::tflite::gpu::metal::SingleOpModel;
@interface ResizeTest : XCTestCase
@end
@implementation ResizeTest
@implementation ResizeTest {
tflite::gpu::metal::MetalExecutionEnvironment exec_env_;
}
- (void)setUp {
[super setUp];
}
@ -195,7 +199,7 @@ using ::tflite::gpu::metal::SingleOpModel;
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeNearestAlignCorners {
- (void)testResizeNearestAlignCornersOp {
TensorRef<BHWC> input;
input.type = DataType::FLOAT32;
input.ref = 0;
@ -221,7 +225,7 @@ using ::tflite::gpu::metal::SingleOpModel;
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeNearestHalfPixelCenters {
- (void)testResizeNearestHalfPixelCentersOp {
TensorRef<BHWC> input;
input.type = DataType::FLOAT32;
input.ref = 0;
@ -247,4 +251,39 @@ using ::tflite::gpu::metal::SingleOpModel;
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeBilinearAligned {
auto status = ResizeBilinearAlignedTest(&exec_env_);
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeBilinearNonAligned {
auto status = ResizeBilinearNonAlignedTest(&exec_env_);
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeBilinearWithoutHalfPixel {
auto status = ResizeBilinearWithoutHalfPixelTest(&exec_env_);
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeBilinearWithHalfPixel {
auto status = ResizeBilinearWithHalfPixelTest(&exec_env_);
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeNearest {
auto status = ResizeNearestTest(&exec_env_);
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeNearestAlignCorners {
auto status = ResizeNearestAlignCornersTest(&exec_env_);
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
- (void)testResizeNearestHalfPixelCenters {
auto status = ResizeNearestHalfPixelCentersTest(&exec_env_);
XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
}
@end