diff --git a/tensorflow/lite/delegates/gpu/cl/kernels/BUILD b/tensorflow/lite/delegates/gpu/cl/kernels/BUILD index eb2dc08f0ef..9786dcce774 100644 --- a/tensorflow/lite/delegates/gpu/cl/kernels/BUILD +++ b/tensorflow/lite/delegates/gpu/cl/kernels/BUILD @@ -266,7 +266,7 @@ cc_test( ":cl_test", "//tensorflow/lite/delegates/gpu/common:operations", "//tensorflow/lite/delegates/gpu/common:status", - "//tensorflow/lite/delegates/gpu/common/tasks:elementwise", + "//tensorflow/lite/delegates/gpu/common/tasks:elementwise_test_util", "@com_google_googletest//:gtest_main", ], ) diff --git a/tensorflow/lite/delegates/gpu/cl/kernels/elementwise_test.cc b/tensorflow/lite/delegates/gpu/cl/kernels/elementwise_test.cc index 9e851f6f162..6fe53a88642 100644 --- a/tensorflow/lite/delegates/gpu/cl/kernels/elementwise_test.cc +++ b/tensorflow/lite/delegates/gpu/cl/kernels/elementwise_test.cc @@ -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/elementwise.h" - #include #include @@ -22,10 +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::FloatEq; -using ::testing::FloatNear; -using ::testing::Pointwise; +#include "tensorflow/lite/delegates/gpu/common/tasks/elementwise_test_util.h" namespace tflite { namespace gpu { @@ -33,1034 +28,183 @@ namespace cl { namespace { TEST_F(OpenCLOperationTest, Abs) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {half(0.0f), half(-1.0f), half(-0.05f), half(0.045f)}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::ABS); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(0.0f), {half(0.0f), half(1.0f), - half(0.05f), half(0.045f)})); - } - } + auto status = AbsTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Cos) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {0.0f, -1.0f, -0.05f, 0.045f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::COS); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT( - dst_tensor.data, - Pointwise(FloatNear(eps), {std::cos(0.0f), std::cos(-1.0f), - std::cos(-0.05f), std::cos(0.045f)})); - } - } + auto status = CosTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Copy) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {half(0.0f), half(-1.0f), half(-0.05f), half(0.045f)}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::COPY); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, Pointwise(FloatEq(), src_tensor.data)); - } - } + auto status = CopyTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Elu) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 1, 1, 7); - src_tensor.data = {0.0f, 1.0f, -1.0f, 100.0f, -100.0f, 0.01f, -0.01f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::ELU); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 1, 1, 7), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.0f, 1.0f, std::exp(-1.0f) - 1.0f, - 100.0f, std::exp(-100.0f) - 1.0f, - 0.01f, std::exp(-0.01f) - 1.0f})); - } - } + auto status = EluTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Exp) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 1, 1, 7); - src_tensor.data = {0.0f, 1.0f, -1.0f, 100.0f, -100.0f, 0.01f, -0.01f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::EXP); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 1, 1, 7), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), - {std::exp(0.0f), std::exp(1.0f), std::exp(-1.0f), - std::exp(100.0f), std::exp(-100.0f), - std::exp(0.01f), std::exp(-0.01f)})); - } - } + auto status = ExpTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, HardSwish) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 1, 1, 7); - src_tensor.data = {-4.5f, -3.0f, -1.5f, 0.0f, 1.5f, 3.0f, 4.5f}; - - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::HARD_SWISH); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - src_tensor.shape, &dst_tensor)); - EXPECT_THAT( - dst_tensor.data, - testing::Pointwise(testing::FloatNear(eps), - {0.0f, 0.0f, -0.375f, 0.0f, 1.125f, 3.f, 4.5f})); - } - } + auto status = HardSwishTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Log) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::LOG); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {std::log(1.0f), std::log(2.0f), - std::log(3.0f), std::log(4.0f)})); - } - } + auto status = LogTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Neg) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {1.0f, -2.0f, 0.0f, 4.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::NEG); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {-1.0f, 2.0f, 0.0f, -4.0f})); - } - } + auto status = NegTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Rsqrt) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::RSQRT); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), - {1.0f / std::sqrt(1.0f), 1.0f / std::sqrt(2.0f), - 1.0f / std::sqrt(3.0f), 1.0f / std::sqrt(4.0f)})); - } - } + auto status = RsqrtTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Sigmoid) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {-std::log(1.0f), -std::log(2.0f), -std::log(3.0f), - -std::log(4.0f)}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::SIGMOID); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.5f, 1.0f / 3.0f, 0.25f, 0.2f})); - } - } + auto status = SigmoidTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Sin) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {0.0f, -1.0f, -0.05f, 0.045f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::SIN); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT( - dst_tensor.data, - Pointwise(FloatNear(eps), {std::sin(0.0f), std::sin(-1.0f), - std::sin(-0.05f), std::sin(0.045f)})); - } - } + auto status = SinTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Sqrt) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::SQRT); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT( - dst_tensor.data, - Pointwise(FloatNear(eps), {std::sqrt(1.0f), std::sqrt(2.0f), - std::sqrt(3.0f), std::sqrt(4.0f)})); - } - } + auto status = SqrtTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Square) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {1.0f, -2.0f, 3.0f, 4.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::SQUARE); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 4.0f, 9.0f, 16.0f})); - } - } + auto status = SquareTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Tanh) { - TensorFloat32 src_tensor; - src_tensor.shape = BHWC(1, 2, 1, 2); - src_tensor.data = {-50.0f, -0.1f, 0.1f, 50.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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; - GPUOperation operation = CreateElementwiseOneInput( - creation_context_.GetGpuInfo(), op_def, OperationType::TANH); - ASSERT_OK(ExecuteGPUOperation( - src_tensor, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT( - dst_tensor.data, - Pointwise(FloatNear(eps), {std::tanh(-50.0f), std::tanh(-0.1f), - std::tanh(0.1f), std::tanh(50.0f)})); - } - } + auto status = TanhTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Sub) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.0f}; - src_tensor_1.data = {0.5f, 1.0f, 3.0f, 3.5f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::SUB, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.5f, 1.0f, 0.0f, 0.5f})); - } - } + auto status = SubTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, SquaredDiff) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.0f}; - src_tensor_1.data = {0.5f, 1.0f, 3.0f, 3.5f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; - 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::SQUARED_DIFF, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.25f, 1.0f, 0.0f, 0.25f})); - } - } + auto status = SquaredDiffTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Div) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; - src_tensor_1.data = {0.5f, 1.0f, 3.0f, 1.5f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::DIV, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {2.0f, 2.0f, 1.0f, 3.0f})); - } - } + auto status = DivTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Pow) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {6.0f, 7.0f, 4.0f, 2.0f}; - src_tensor_1.data = {0.0f, 1.0f, 2.0f, 3.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::POW, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 7.0f, 16.0f, 8.0f})); - } - } + auto status = PowTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Add) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; - src_tensor_1.data = {0.5f, 1.0f, 3.0f, 1.5f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::ADD, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.5f, 3.0f, 6.0f, 6.0f})); - } - } + auto status = AddTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Maximum) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; - src_tensor_1.data = {1.0f, 2.0f, 3.0f, -2.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::MAXIMUM, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 2.0f, 3.0f, -2.0f})); - } - } + auto status = MaximumTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, MaximumWithScalar) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 4, 1, 1); - src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; - - ElementwiseAttributes attr; - attr.param = -1.0f; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwise( - creation_context_.GetGpuInfo(), op_def, OperationType::MAXIMUM, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 4, 1, 1), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.0f, -1.0f, 2.0f, -1.0f})); - } - } + auto status = MaximumWithScalarTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, MaximumWithConstantLinearTensor) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {1.0f, -6.2f, -2.0f, 3.0f}; - - ::tflite::gpu::Tensor linear_tensor; - linear_tensor.shape = Linear(2); - linear_tensor.data = {0.5f, 2.0f}; - ElementwiseAttributes attr; - attr.param = linear_tensor; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwise( - creation_context_.GetGpuInfo(), op_def, OperationType::MAXIMUM, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 2.0f, 0.5f, 3.0f})); - } - } + auto status = MaximumWithConstantLinearTensorTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, MaximumWithConstantHWCTensor) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {1.0f, -6.2f, -2.0f, 3.0f}; - - ::tflite::gpu::Tensor hwc_tensor; - hwc_tensor.shape = HWC(2, 1, 2); - hwc_tensor.data = {0.5f, 2.0f, 0.7f, 4.7f}; - ElementwiseAttributes attr; - attr.param = hwc_tensor; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwise( - creation_context_.GetGpuInfo(), op_def, OperationType::MAXIMUM, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 2.0f, 0.7f, 4.7f})); - } - } + auto status = MaximumWithConstantHWCTensorTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } + TEST_F(OpenCLOperationTest, MaximumWithConstantHWCTensorBroadcastChannels) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {1.0f, -6.2f, -2.0f, 3.0f}; - - ::tflite::gpu::Tensor hwc_tensor; - hwc_tensor.shape = HWC(2, 1, 1); - hwc_tensor.data = {0.5f, 2.0f}; - ElementwiseAttributes attr; - attr.param = hwc_tensor; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwise( - creation_context_.GetGpuInfo(), op_def, OperationType::MAXIMUM, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 0.5f, 2.0f, 3.0f})); - } - } + auto status = MaximumWithConstantHWCTensorBroadcastChannelsTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Minimum) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; - src_tensor_1.data = {1.0f, 2.0f, 3.0f, -2.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::MINIMUM, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.0f, -6.2f, 2.0f, -3.0f})); - } - } + auto status = MinimumTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, MinimumWithScalar) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 4, 1, 1); - src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; - - ElementwiseAttributes attr; - attr.param = -1.0f; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwise( - creation_context_.GetGpuInfo(), op_def, OperationType::MINIMUM, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 4, 1, 1), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {-1.0f, -6.2f, -1.0f, -3.0f})); - } - } + auto status = MinimumWithScalarTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Mul) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; - src_tensor_1.data = {0.5f, 1.0f, 3.0f, 1.5f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::MUL, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.5f, 2.0f, 9.0f, 6.75f})); - } - } + auto status = MulTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, MulBroadcastHW) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 1, 1, 2); - src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; - src_tensor_1.data = {0.5f, 3.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::MUL, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.5f, 6.0f, 1.5f, 13.5f})); - } - } + auto status = MulBroadcastHWTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, MulBroadcastChannels) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 1); - src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; - src_tensor_1.data = {0.5f, 3.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::MUL, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.5f, 1.0f, 9.0f, 13.5f})); - } - } + auto status = MulBroadcastChannelsTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, SubWithScalarAtFirstPosition) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 4, 1, 1); - src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; - - ElementwiseAttributes attr; - attr.param = 4.0f; - attr.runtime_tensor_is_second = true; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwise( - creation_context_.GetGpuInfo(), op_def, OperationType::SUB, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 4, 1, 1), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {4.0f, 10.2f, 2.0f, 7.0f})); - } - } + auto status = SubWithScalarAtFirstPositionTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Less) { - TensorFloat32 src_tensor_0, src_tensor_1; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_1.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; - src_tensor_1.data = {1.0f, 0.0f, 2.0f, -4.0f}; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); - op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); - TensorFloat32 dst_tensor; - GPUOperation operation = CreateElementwiseTwoInput( - op_def, OperationType::LESS, src_tensor_1.shape); - ASSERT_OK(ExecuteGPUOperation( - {src_tensor_0, src_tensor_1}, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 0.0f, 0.0f, 0.0f})); - } - } + auto status = LessTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, LessEqual) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; - - ElementwiseAttributes attr; - attr.param = 2.0f; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = - CreateElementwise(creation_context_.GetGpuInfo(), op_def, - OperationType::LESS_EQUAL, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 1.0f, 1.0f, 0.0f})); - } - } + auto status = LessEqualTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Greater) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; - - ElementwiseAttributes attr; - attr.param = 2.0f; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwise( - creation_context_.GetGpuInfo(), op_def, OperationType::GREATER, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.0f, 0.0f, 0.0f, 1.0f})); - } - } + auto status = GreaterTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, GreaterEqual) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; - - ElementwiseAttributes attr; - attr.param = 2.0f; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = - CreateElementwise(creation_context_.GetGpuInfo(), op_def, - OperationType::GREATER_EQUAL, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.0f, 0.0f, 1.0f, 1.0f})); - } - } + auto status = GreaterEqualTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, Equal) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; - - ElementwiseAttributes attr; - attr.param = 2.0f; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = CreateElementwise( - creation_context_.GetGpuInfo(), op_def, OperationType::EQUAL, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {0.0f, 0.0f, 1.0f, 0.0f})); - } - } + auto status = EqualTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } TEST_F(OpenCLOperationTest, NotEqual) { - TensorFloat32 src_tensor_0; - src_tensor_0.shape = BHWC(1, 2, 1, 2); - src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; - - ElementwiseAttributes attr; - attr.param = 2.0f; - - for (auto storage : env_.GetSupportedStorages()) { - for (auto precision : env_.GetSupportedPrecisions()) { - const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; - GPUOperation operation = - CreateElementwise(creation_context_.GetGpuInfo(), op_def, - OperationType::NOT_EQUAL, attr); - ASSERT_OK(ExecuteGPUOperation( - src_tensor_0, creation_context_, - absl::make_unique(std::move(operation)), - BHWC(1, 2, 1, 2), &dst_tensor)); - EXPECT_THAT(dst_tensor.data, - Pointwise(FloatNear(eps), {1.0f, 1.0f, 0.0f, 1.0f})); - } - } + auto status = NotEqualTest(&exec_env_); + ASSERT_TRUE(status.ok()) << status.error_message(); } } // namespace diff --git a/tensorflow/lite/delegates/gpu/common/tasks/BUILD b/tensorflow/lite/delegates/gpu/common/tasks/BUILD index b7ac860759e..27a7a059f35 100644 --- a/tensorflow/lite/delegates/gpu/common/tasks/BUILD +++ b/tensorflow/lite/delegates/gpu/common/tasks/BUILD @@ -295,6 +295,19 @@ cc_library( ], ) +cc_library( + name = "elementwise_test_util", + testonly = 1, + srcs = ["elementwise_test_util.cc"], + hdrs = ["elementwise_test_util.h"], + deps = [ + ":elementwise", + "//tensorflow/lite/delegates/gpu/common:operations", + "//tensorflow/lite/delegates/gpu/common:status", + "//tensorflow/lite/delegates/gpu/common/task:testing_util", + ], +) + cc_library( name = "fully_connected", srcs = ["fully_connected.cc"], diff --git a/tensorflow/lite/delegates/gpu/common/tasks/elementwise_test_util.cc b/tensorflow/lite/delegates/gpu/common/tasks/elementwise_test_util.cc new file mode 100644 index 00000000000..bd77bf3cee5 --- /dev/null +++ b/tensorflow/lite/delegates/gpu/common/tasks/elementwise_test_util.cc @@ -0,0 +1,1064 @@ +/* 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/elementwise_test_util.h" + +#include + +#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/elementwise.h" + +namespace tflite { +namespace gpu { + +absl::Status AbsTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {half(0.0f), half(-1.0f), half(-0.05f), half(0.045f)}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::ABS); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({half(0.0f), half(1.0f), half(0.05f), half(0.045f)}, + dst_tensor.data, 0.0f)); + } + } + return absl::OkStatus(); +} + +absl::Status CosTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {0.0f, -1.0f, -0.05f, 0.045f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 5e-5f : 1e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::COS); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear( + {std::cos(0.0f), std::cos(-1.0f), std::cos(-0.05f), std::cos(0.045f)}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status CopyTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {half(0.0f), half(-1.0f), half(-0.05f), half(0.045f)}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::COPY); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear(src_tensor.data, dst_tensor.data, 0.0f)); + } + } + return absl::OkStatus(); +} + +absl::Status EluTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 1, 1, 7); + src_tensor.data = {0.0f, 1.0f, -1.0f, 100.0f, -100.0f, 0.01f, -0.01f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::ELU); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 1, 1, 7), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear( + {0.0f, 1.0f, std::exp(-1.0f) - 1.0f, 100.0f, std::exp(-100.0f) - 1.0f, + 0.01f, std::exp(-0.01f) - 1.0f}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status ExpTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 1, 1, 7); + src_tensor.data = {0.0f, 1.0f, -1.0f, 100.0f, -100.0f, 0.01f, -0.01f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::EXP); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 1, 1, 7), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear( + {std::exp(0.0f), std::exp(1.0f), std::exp(-1.0f), std::exp(100.0f), + std::exp(-100.0f), std::exp(0.01f), std::exp(-0.01f)}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status HardSwishTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 1, 1, 7); + src_tensor.data = {-4.5f, -3.0f, -1.5f, 0.0f, 1.5f, 3.0f, 4.5f}; + + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::HARD_SWISH); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + src_tensor.shape, &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.0f, 0.0f, -0.375f, 0.0f, 1.125f, 3.f, 4.5f}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status LogTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::LOG); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear( + {std::log(1.0f), std::log(2.0f), std::log(3.0f), std::log(4.0f)}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status NegTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {1.0f, -2.0f, 0.0f, 4.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::NEG); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({-1.0f, 2.0f, 0.0f, -4.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status RsqrtTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::RSQRT); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f / std::sqrt(1.0f), 1.0f / std::sqrt(2.0f), + 1.0f / std::sqrt(3.0f), 1.0f / std::sqrt(4.0f)}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status SigmoidTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {-std::log(1.0f), -std::log(2.0f), -std::log(3.0f), + -std::log(4.0f)}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::SIGMOID); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear({0.5f, 1.0f / 3.0f, 0.25f, 0.2f}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status SinTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {0.0f, -1.0f, -0.05f, 0.045f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-5f : 5e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::SIN); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear( + {std::sin(0.0f), std::sin(-1.0f), std::sin(-0.05f), std::sin(0.045f)}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status SqrtTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {1.0f, 2.0f, 3.0f, 4.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::SQRT); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear( + {std::sqrt(1.0f), std::sqrt(2.0f), std::sqrt(3.0f), std::sqrt(4.0f)}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status SquareTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {1.0f, -2.0f, 3.0f, 4.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::SQUARE); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 4.0f, 9.0f, 16.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status TanhTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor; + src_tensor.shape = BHWC(1, 2, 1, 2); + src_tensor.data = {-50.0f, -0.1f, 0.1f, 50.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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; + GPUOperation operation = CreateElementwiseOneInput( + env->GetGpuInfo(), op_def, OperationType::TANH); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR(PointWiseNear({std::tanh(-50.0f), std::tanh(-0.1f), + std::tanh(0.1f), std::tanh(50.0f)}, + dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status SubTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.0f}; + src_tensor_1.data = {0.5f, 1.0f, 3.0f, 3.5f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::SUB, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.5f, 1.0f, 0.0f, 0.5f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status SquaredDiffTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.0f}; + src_tensor_1.data = {0.5f, 1.0f, 3.0f, 3.5f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 1e-3f; + 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::SQUARED_DIFF, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.25f, 1.0f, 0.0f, 0.25f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status DivTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; + src_tensor_1.data = {0.5f, 1.0f, 3.0f, 1.5f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::DIV, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({2.0f, 2.0f, 1.0f, 3.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status PowTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {6.0f, 7.0f, 4.0f, 2.0f}; + src_tensor_1.data = {0.0f, 1.0f, 2.0f, 3.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::POW, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 7.0f, 16.0f, 8.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status AddTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; + src_tensor_1.data = {0.5f, 1.0f, 3.0f, 1.5f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::ADD, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.5f, 3.0f, 6.0f, 6.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MaximumTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; + src_tensor_1.data = {1.0f, 2.0f, 3.0f, -2.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::MAXIMUM, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 2.0f, 3.0f, -2.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MaximumWithScalarTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 4, 1, 1); + src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; + + ElementwiseAttributes attr; + attr.param = -1.0f; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise(env->GetGpuInfo(), op_def, + OperationType::MAXIMUM, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 4, 1, 1), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.0f, -1.0f, 2.0f, -1.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MaximumWithConstantLinearTensorTest( + TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {1.0f, -6.2f, -2.0f, 3.0f}; + + ::tflite::gpu::Tensor linear_tensor; + linear_tensor.shape = Linear(2); + linear_tensor.data = {0.5f, 2.0f}; + ElementwiseAttributes attr; + attr.param = linear_tensor; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise(env->GetGpuInfo(), op_def, + OperationType::MAXIMUM, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 2.0f, 0.5f, 3.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MaximumWithConstantHWCTensorTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {1.0f, -6.2f, -2.0f, 3.0f}; + + ::tflite::gpu::Tensor hwc_tensor; + hwc_tensor.shape = HWC(2, 1, 2); + hwc_tensor.data = {0.5f, 2.0f, 0.7f, 4.7f}; + ElementwiseAttributes attr; + attr.param = hwc_tensor; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise(env->GetGpuInfo(), op_def, + OperationType::MAXIMUM, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 2.0f, 0.7f, 4.7f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} +absl::Status MaximumWithConstantHWCTensorBroadcastChannelsTest( + TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {1.0f, -6.2f, -2.0f, 3.0f}; + + ::tflite::gpu::Tensor hwc_tensor; + hwc_tensor.shape = HWC(2, 1, 1); + hwc_tensor.data = {0.5f, 2.0f}; + ElementwiseAttributes attr; + attr.param = hwc_tensor; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise(env->GetGpuInfo(), op_def, + OperationType::MAXIMUM, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 0.5f, 2.0f, 3.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MinimumTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; + src_tensor_1.data = {1.0f, 2.0f, 3.0f, -2.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::MINIMUM, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.0f, -6.2f, 2.0f, -3.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MinimumWithScalarTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 4, 1, 1); + src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; + + ElementwiseAttributes attr; + attr.param = -1.0f; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise(env->GetGpuInfo(), op_def, + OperationType::MINIMUM, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 4, 1, 1), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({-1.0f, -6.2f, -1.0f, -3.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MulTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; + src_tensor_1.data = {0.5f, 1.0f, 3.0f, 1.5f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::MUL, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.5f, 2.0f, 9.0f, 6.75f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MulBroadcastHWTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 1, 1, 2); + src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; + src_tensor_1.data = {0.5f, 3.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::MUL, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.5f, 6.0f, 1.5f, 13.5f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status MulBroadcastChannelsTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 1); + src_tensor_0.data = {1.0f, 2.0f, 3.0f, 4.5f}; + src_tensor_1.data = {0.5f, 3.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::MUL, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.5f, 1.0f, 9.0f, 13.5f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status SubWithScalarAtFirstPositionTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 4, 1, 1); + src_tensor_0.data = {0.0f, -6.2f, 2.0f, -3.0f}; + + ElementwiseAttributes attr; + attr.param = 4.0f; + attr.runtime_tensor_is_second = true; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise(env->GetGpuInfo(), op_def, + OperationType::SUB, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 4, 1, 1), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({4.0f, 10.2f, 2.0f, 7.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status LessTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0, src_tensor_1; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_1.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; + src_tensor_1.data = {1.0f, 0.0f, 2.0f, -4.0f}; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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.src_tensors.push_back({data_type, storage, Layout::HWC}); + op_def.dst_tensors.push_back({data_type, storage, Layout::HWC}); + TensorFloat32 dst_tensor; + GPUOperation operation = CreateElementwiseTwoInput( + op_def, OperationType::LESS, src_tensor_1.shape); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + {src_tensor_0, src_tensor_1}, + absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 0.0f, 0.0f, 0.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status LessEqualTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; + + ElementwiseAttributes attr; + attr.param = 2.0f; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise( + env->GetGpuInfo(), op_def, OperationType::LESS_EQUAL, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 1.0f, 1.0f, 0.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status GreaterTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; + + ElementwiseAttributes attr; + attr.param = 2.0f; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise(env->GetGpuInfo(), op_def, + OperationType::GREATER, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.0f, 0.0f, 0.0f, 1.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status GreaterEqualTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; + + ElementwiseAttributes attr; + attr.param = 2.0f; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise( + env->GetGpuInfo(), op_def, OperationType::GREATER_EQUAL, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.0f, 0.0f, 1.0f, 1.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status EqualTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; + + ElementwiseAttributes attr; + attr.param = 2.0f; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise(env->GetGpuInfo(), op_def, + OperationType::EQUAL, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({0.0f, 0.0f, 1.0f, 0.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +absl::Status NotEqualTest(TestExecutionEnvironment* env) { + TensorFloat32 src_tensor_0; + src_tensor_0.shape = BHWC(1, 2, 1, 2); + src_tensor_0.data = {0.0f, 1.0f, 2.0f, 3.0f}; + + ElementwiseAttributes attr; + attr.param = 2.0f; + + for (auto storage : env->GetSupportedStorages()) { + for (auto precision : env->GetSupportedPrecisions()) { + const float eps = precision == CalculationsPrecision::F32 ? 1e-6f : 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; + GPUOperation operation = CreateElementwise( + env->GetGpuInfo(), op_def, OperationType::NOT_EQUAL, attr); + RETURN_IF_ERROR(env->ExecuteGPUOperation( + src_tensor_0, absl::make_unique(std::move(operation)), + BHWC(1, 2, 1, 2), &dst_tensor)); + RETURN_IF_ERROR( + PointWiseNear({1.0f, 1.0f, 0.0f, 1.0f}, dst_tensor.data, eps)); + } + } + return absl::OkStatus(); +} + +} // namespace gpu +} // namespace tflite diff --git a/tensorflow/lite/delegates/gpu/common/tasks/elementwise_test_util.h b/tensorflow/lite/delegates/gpu/common/tasks/elementwise_test_util.h new file mode 100644 index 00000000000..3e307244048 --- /dev/null +++ b/tensorflow/lite/delegates/gpu/common/tasks/elementwise_test_util.h @@ -0,0 +1,66 @@ +/* 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_ELEMENTWISE_TEST_UTIL_H_ +#define TENSORFLOW_LITE_DELEGATES_GPU_COMMON_TASKS_ELEMENTWISE_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 AbsTest(TestExecutionEnvironment* env); +absl::Status CosTest(TestExecutionEnvironment* env); +absl::Status CopyTest(TestExecutionEnvironment* env); +absl::Status EluTest(TestExecutionEnvironment* env); +absl::Status ExpTest(TestExecutionEnvironment* env); +absl::Status HardSwishTest(TestExecutionEnvironment* env); +absl::Status LogTest(TestExecutionEnvironment* env); +absl::Status NegTest(TestExecutionEnvironment* env); +absl::Status RsqrtTest(TestExecutionEnvironment* env); +absl::Status SigmoidTest(TestExecutionEnvironment* env); +absl::Status SinTest(TestExecutionEnvironment* env); +absl::Status SqrtTest(TestExecutionEnvironment* env); +absl::Status SquareTest(TestExecutionEnvironment* env); +absl::Status TanhTest(TestExecutionEnvironment* env); +absl::Status SubTest(TestExecutionEnvironment* env); +absl::Status SquaredDiffTest(TestExecutionEnvironment* env); +absl::Status DivTest(TestExecutionEnvironment* env); +absl::Status PowTest(TestExecutionEnvironment* env); +absl::Status AddTest(TestExecutionEnvironment* env); +absl::Status MaximumTest(TestExecutionEnvironment* env); +absl::Status MaximumWithScalarTest(TestExecutionEnvironment* env); +absl::Status MaximumWithConstantLinearTensorTest(TestExecutionEnvironment* env); +absl::Status MaximumWithConstantHWCTensorTest(TestExecutionEnvironment* env); +absl::Status MaximumWithConstantHWCTensorBroadcastChannelsTest( + TestExecutionEnvironment* env); +absl::Status MinimumTest(TestExecutionEnvironment* env); +absl::Status MinimumWithScalarTest(TestExecutionEnvironment* env); +absl::Status MulTest(TestExecutionEnvironment* env); +absl::Status MulBroadcastHWTest(TestExecutionEnvironment* env); +absl::Status MulBroadcastChannelsTest(TestExecutionEnvironment* env); +absl::Status SubWithScalarAtFirstPositionTest(TestExecutionEnvironment* env); +absl::Status LessTest(TestExecutionEnvironment* env); +absl::Status LessEqualTest(TestExecutionEnvironment* env); +absl::Status GreaterTest(TestExecutionEnvironment* env); +absl::Status GreaterEqualTest(TestExecutionEnvironment* env); +absl::Status EqualTest(TestExecutionEnvironment* env); +absl::Status NotEqualTest(TestExecutionEnvironment* env); + +} // namespace gpu +} // namespace tflite + +#endif // TENSORFLOW_LITE_DELEGATES_GPU_COMMON_TASKS_ELEMENTWISE_TEST_UTIL_H_ diff --git a/tensorflow/lite/delegates/gpu/metal/kernels/BUILD b/tensorflow/lite/delegates/gpu/metal/kernels/BUILD index e969e364a77..c54be6ad4f3 100644 --- a/tensorflow/lite/delegates/gpu/metal/kernels/BUILD +++ b/tensorflow/lite/delegates/gpu/metal/kernels/BUILD @@ -231,6 +231,7 @@ objc_library( deps = [ ":elementwise", ":test_util", + "//tensorflow/lite/delegates/gpu/common/tasks:elementwise_test_util", ], ) @@ -904,6 +905,7 @@ objc_library( "//tensorflow/lite/delegates/gpu/common:types", "//tensorflow/lite/delegates/gpu/common:util", "//tensorflow/lite/delegates/gpu/common/tasks:add_test_util", + "//tensorflow/lite/delegates/gpu/common/tasks:elementwise_test_util", "//tensorflow/lite/delegates/gpu/common/tasks:prelu_test_util", "//tensorflow/lite/delegates/gpu/common/tasks:quantize_and_dequantize_test_util", "//tensorflow/lite/delegates/gpu/common/tasks:relu_test_util", diff --git a/tensorflow/lite/delegates/gpu/metal/kernels/elementwise_test.mm b/tensorflow/lite/delegates/gpu/metal/kernels/elementwise_test.mm index 03dff9b83ce..b745d633061 100644 --- a/tensorflow/lite/delegates/gpu/metal/kernels/elementwise_test.mm +++ b/tensorflow/lite/delegates/gpu/metal/kernels/elementwise_test.mm @@ -13,7 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. ==============================================================================*/ -#include "tensorflow/lite/delegates/gpu/metal/kernels/add.h" +#include "tensorflow/lite/delegates/gpu/metal/kernels/elementwise.h" #import @@ -27,6 +27,7 @@ limitations under the License. #include "tensorflow/lite/delegates/gpu/common/util.h" #include "tensorflow/lite/delegates/gpu/metal/compute_task_descriptor.h" #include "tensorflow/lite/delegates/gpu/metal/kernels/test_util.h" +#include "tensorflow/lite/delegates/gpu/common/tasks/elementwise_test_util.h" using ::tflite::gpu::DataType; using ::tflite::gpu::HWC; @@ -39,7 +40,10 @@ using ::tflite::gpu::metal::SingleOpModel; @interface ElementwiseTest : XCTestCase @end -@implementation ElementwiseTest +@implementation ElementwiseTest { + tflite::gpu::metal::MetalExecutionEnvironment exec_env_; +} + - (void)setUp { [super setUp]; } @@ -419,4 +423,184 @@ TensorRef GetTensorRef(int ref, const BHWC& shape) { XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); } +- (void)testAbsUnit { + auto status = AbsTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testCosUnit { + auto status = CosTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testCopyUnit { + auto status = CopyTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testEluUnit { + auto status = EluTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testExpUnit { + auto status = ExpTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testHardSwishUnit { + auto status = HardSwishTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testLogUnit { + auto status = LogTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testNegUnit { + auto status = NegTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testRsqrtUnit { + auto status = RsqrtTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testSigmoidUnit { + auto status = SigmoidTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testSinUnit { + auto status = SinTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testSqrtUnit { + auto status = SqrtTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testSquareUnit { + auto status = SquareTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testTanhUnit { + auto status = TanhTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testSubUnit { + auto status = SubTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testSquaredDiffUnit { + auto status = SquaredDiffTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testDivUnit { + auto status = DivTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testPowUnit { + auto status = PowTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testAddUnit { + auto status = AddTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMaximumUnit { + auto status = MaximumTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMaximumWithScalarUnit { + auto status = MaximumWithScalarTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMaximumWithConstantLinearTensorUnit { + auto status = MaximumWithConstantLinearTensorTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMaximumWithConstantHWCTensorUnit { + auto status = MaximumWithConstantHWCTensorTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMaximumWithConstantHWCTensorBroadcastChannelsUnit { + auto status = MaximumWithConstantHWCTensorBroadcastChannelsTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMinimumUnit { + auto status = MinimumTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMinimumWithScalarUnit { + auto status = MinimumWithScalarTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMulUnit { + auto status = MulTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMulBroadcastHWUnit { + auto status = MulBroadcastHWTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testMulBroadcastChannelsUnit { + auto status = MulBroadcastChannelsTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testSubWithScalarAtFirstPositionUnit { + auto status = SubWithScalarAtFirstPositionTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testLessUnit { + auto status = LessTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testLessEqualUnit { + auto status = LessEqualTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testGreaterUnit { + auto status = GreaterTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testGreaterEqualUnit { + auto status = GreaterEqualTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testEqualUnit { + auto status = EqualTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + +- (void)testNotEqualUnit { + auto status = NotEqualTest(&exec_env_); + XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str()); +} + @end