Added ConverterToConvWeights unit tests for OpenCL & Metal.

PiperOrigin-RevId: 356489082
Change-Id: I9d0aa4e6ddf084744e7cba5466d25a4cb761c142

tensorflow/lite/delegates/gpu/cl/kernels/BUILD

@@ -111,6 +111,23 @@ cc_test(
     ],
 )
 
+cc_test(
+    name = "conv_weights_converter_test",
+    srcs = ["conv_weights_converter_test.cc"],
+    linkstatic = True,
+    tags = tf_gpu_tests_tags() + [
+        "linux",
+        "local",
+    ],
+    deps = [
+        ":cl_test",
+        "//tensorflow/lite/delegates/gpu/common:operations",
+        "//tensorflow/lite/delegates/gpu/common:status",
+        "//tensorflow/lite/delegates/gpu/common/tasks:conv_weights_converter_test_util",
+        "@com_google_googletest//:gtest_main",
+    ],
+)
+
 cc_library(
     name = "converter",
     srcs = ["converter.cc"],
@@ -627,6 +644,7 @@ test_suite(
         "conv_buffer_1x1_test",
         "conv_constants_test",
         "conv_powervr_test",
+        "conv_weights_converter_test",
         "convolution_transposed_3x3_test",
         "convolution_transposed_3x3_thin_test",
         "convolution_transposed_4x4_test",

tensorflow/lite/delegates/gpu/cl/kernels/conv_weights_converter_test.cc

@@ -0,0 +1,56 @@
+/* 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 <vector>
+
+#include <gmock/gmock.h>
+#include <gtest/gtest.h>
+#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"
+#include "tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter_test_util.h"
+
+namespace tflite {
+namespace gpu {
+namespace cl {
+
+TEST_F(OpenCLOperationTest, ConverterToConvWeights1x1OutX4) {
+  const auto status = ConverterToConvWeights1x1OutX4Test(&exec_env_);
+  ASSERT_TRUE(status.ok()) << status.error_message();
+}
+
+TEST_F(OpenCLOperationTest, ConverterToConvWeights1x1OutX4Unaligned) {
+  const auto status = ConverterToConvWeights1x1OutX4UnalignedTest(&exec_env_);
+  ASSERT_TRUE(status.ok()) << status.error_message();
+}
+
+TEST_F(OpenCLOperationTest, ConverterToConvWeights1x1OutX2) {
+  const auto status = ConverterToConvWeights1x1OutX2Test(&exec_env_);
+  ASSERT_TRUE(status.ok()) << status.error_message();
+}
+
+TEST_F(OpenCLOperationTest, ConverterToConvWeightsOutX2) {
+  const auto status = ConverterToConvWeightsOutX2Test(&exec_env_);
+  ASSERT_TRUE(status.ok()) << status.error_message();
+}
+
+TEST_F(OpenCLOperationTest, ConverterToConvTransposedWeights4x4) {
+  const auto status = ConverterToConvTransposedWeights4x4Test(&exec_env_);
+  ASSERT_TRUE(status.ok()) << status.error_message();
+}
+
+}  // namespace cl
+}  // namespace gpu
+}  // namespace tflite

tensorflow/lite/delegates/gpu/common/tasks/BUILD

@@ -150,6 +150,20 @@ cc_library(
     ],
 )
 
+cc_library(
+    name = "conv_weights_converter_test_util",
+    testonly = 1,
+    srcs = ["conv_weights_converter_test_util.cc"],
+    hdrs = ["conv_weights_converter_test_util.h"],
+    deps = [
+        ":conv_weights_converter",
+        "//tensorflow/lite/delegates/gpu/common:operations",
+        "//tensorflow/lite/delegates/gpu/common:status",
+        "//tensorflow/lite/delegates/gpu/common/task:testing_util",
+        "//tensorflow/lite/delegates/gpu/common/task:weights_conversion",
+    ],
+)
+
 cc_library(
     name = "convolution_transposed",
     srcs = ["convolution_transposed.cc"],

tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter_test_util.cc

@@ -0,0 +1,242 @@
+/* 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/conv_weights_converter_test_util.h"
+
+#include <vector>
+
+#include "tensorflow/lite/delegates/gpu/common/operations.h"
+#include "tensorflow/lite/delegates/gpu/common/status.h"
+#include "tensorflow/lite/delegates/gpu/common/task/testing_util.h"
+#include "tensorflow/lite/delegates/gpu/common/task/weights_conversion.h"
+#include "tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter.h"
+
+namespace tflite {
+namespace gpu {
+namespace {
+absl::Status ConvolutionWeightsConverterTest(
+    const Tensor<OHWI, DataType::FLOAT32>& weights,
+    const WeightsDescription& weight_desc, TestExecutionEnvironment* env,
+    const OperationDef& op_def) {
+  TensorFloat32 dst_tensor;
+  const int flt_count =
+      GetTotalElementsCountForLayout(weight_desc, weights.shape);
+  dst_tensor.shape = BHWC(1, 1, 1, flt_count);
+  dst_tensor.data.resize(flt_count);
+  RearrangeWeights(
+      weights, weight_desc, DataType::FLOAT32,
+      absl::MakeSpan(reinterpret_cast<uint8_t*>(dst_tensor.data.data()),
+                     flt_count * 4));
+
+  // reinterpreting weights in OHWI as tensor in BHWC
+  TensorFloat32 src_tensor;
+  auto src_shape =
+      BHWC(weights.shape.o, weights.shape.h, weights.shape.w, weights.shape.i);
+  src_tensor.shape = src_shape;
+  src_tensor.data.resize(src_shape.DimensionsProduct(), 2.0);
+  for (int o = 0; o < weights.shape.o; ++o) {
+    for (int y = 0; y < weights.shape.h; ++y) {
+      for (int x = 0; x < weights.shape.w; ++x) {
+        for (int i = 0; i < weights.shape.i; ++i) {
+          const int f_index = weights.shape.LinearIndex({o, y, x, i});
+          const int s_index = src_shape.LinearIndex({o, y, x, i});
+          src_tensor.data[s_index] = weights.data[f_index];
+        }
+      }
+    }
+  }
+
+  TensorFloat32 dst_tensor_gpu;
+  auto converter = ConverterToConvWeights(op_def, weight_desc);
+  RETURN_IF_ERROR(env->ExecuteGPUOperation(
+      src_tensor,
+      absl::make_unique<ConverterToConvWeights>(std::move(converter)),
+      dst_tensor.shape, &dst_tensor_gpu));
+  RETURN_IF_ERROR(PointWiseNear(dst_tensor.data, dst_tensor_gpu.data, 0.0f));
+  return absl::OkStatus();
+}
+
+}  // namespace
+
+absl::Status ConverterToConvWeights1x1OutX4Test(TestExecutionEnvironment* env) {
+  const int kSrcChannels = 8;
+  const int kDstChannels = 32;
+  auto weights_shape = OHWI(kDstChannels, 1, 1, kSrcChannels);
+  WeightsDescription conv_weight_desc;
+  conv_weight_desc.output_group_size = 4;
+
+  Tensor<OHWI, DataType::FLOAT32> weights;
+  weights.shape = weights_shape;
+  weights.data.resize(weights_shape.DimensionsProduct());
+  for (int i = 0; i < weights.data.size(); ++i) {
+    weights.data[i] = half(static_cast<float>(i));
+  }
+
+  for (auto storage : env->GetSupportedStorages()) {
+    for (auto precision : env->GetSupportedPrecisions()) {
+      for (auto weights_layout :
+           {WeightsLayout::kOHWIOGroupI4O4, WeightsLayout::kOHWIOGroupO4I4}) {
+        conv_weight_desc.layout = weights_layout;
+        OperationDef op_def;
+        op_def.precision = precision;
+        auto data_type = DeduceDataTypeFromPrecision(precision);
+        op_def.src_tensors.push_back({data_type, storage, Layout::BHWC});
+        op_def.dst_tensors.push_back(
+            {data_type, TensorStorageType::BUFFER, Layout::UNKNOWN});
+        RETURN_IF_ERROR(ConvolutionWeightsConverterTest(
+            weights, conv_weight_desc, env, op_def));
+      }
+    }
+  }
+  return absl::OkStatus();
+}
+
+absl::Status ConverterToConvWeights1x1OutX4UnalignedTest(
+    TestExecutionEnvironment* env) {
+  const int kSrcChannels = 8;
+  const int kDstChannels = 17;
+  auto weights_shape = OHWI(kDstChannels, 1, 1, kSrcChannels);
+  WeightsDescription conv_weight_desc;
+  conv_weight_desc.output_group_size = 4;
+
+  Tensor<OHWI, DataType::FLOAT32> weights;
+  weights.shape = weights_shape;
+  weights.data.resize(weights_shape.DimensionsProduct());
+  for (int i = 0; i < weights.data.size(); ++i) {
+    weights.data[i] = half(static_cast<float>(i));
+  }
+
+  for (auto storage : env->GetSupportedStorages()) {
+    for (auto precision : env->GetSupportedPrecisions()) {
+      for (auto weights_layout :
+           {WeightsLayout::kOHWIOGroupI4O4, WeightsLayout::kOHWIOGroupO4I4}) {
+        conv_weight_desc.layout = weights_layout;
+        OperationDef op_def;
+        op_def.precision = precision;
+        auto data_type = DeduceDataTypeFromPrecision(precision);
+        op_def.src_tensors.push_back({data_type, storage, Layout::BHWC});
+        op_def.dst_tensors.push_back(
+            {data_type, TensorStorageType::BUFFER, Layout::UNKNOWN});
+        RETURN_IF_ERROR(ConvolutionWeightsConverterTest(
+            weights, conv_weight_desc, env, op_def));
+      }
+    }
+  }
+  return absl::OkStatus();
+}
+
+absl::Status ConverterToConvWeights1x1OutX2Test(TestExecutionEnvironment* env) {
+  const int kSrcChannels = 7;
+  const int kDstChannels = 37;
+  auto weights_shape = OHWI(kDstChannels, 1, 1, kSrcChannels);
+  WeightsDescription conv_weight_desc;
+  conv_weight_desc.output_group_size = 2;
+
+  Tensor<OHWI, DataType::FLOAT32> weights;
+  weights.shape = weights_shape;
+  weights.data.resize(weights_shape.DimensionsProduct());
+  for (int i = 0; i < weights.data.size(); ++i) {
+    weights.data[i] = half(static_cast<float>(i));
+  }
+
+  for (auto storage : env->GetSupportedStorages()) {
+    for (auto precision : env->GetSupportedPrecisions()) {
+      for (auto weights_layout :
+           {WeightsLayout::kOHWIOGroupI4O4, WeightsLayout::kOHWIOGroupO4I4}) {
+        conv_weight_desc.layout = weights_layout;
+        OperationDef op_def;
+        op_def.precision = precision;
+        auto data_type = DeduceDataTypeFromPrecision(precision);
+        op_def.src_tensors.push_back({data_type, storage, Layout::BHWC});
+        op_def.dst_tensors.push_back(
+            {data_type, TensorStorageType::BUFFER, Layout::UNKNOWN});
+        RETURN_IF_ERROR(ConvolutionWeightsConverterTest(
+            weights, conv_weight_desc, env, op_def));
+      }
+    }
+  }
+  return absl::OkStatus();
+}
+
+absl::Status ConverterToConvWeightsOutX2Test(TestExecutionEnvironment* env) {
+  const int kSrcChannels = 8;
+  const int kDstChannels = 38;
+  auto weights_shape = OHWI(kDstChannels, 3, 4, kSrcChannels);
+  WeightsDescription conv_weight_desc;
+  conv_weight_desc.output_group_size = 2;
+
+  Tensor<OHWI, DataType::FLOAT32> weights;
+  weights.shape = weights_shape;
+  weights.data.resize(weights_shape.DimensionsProduct());
+  for (int i = 0; i < weights.data.size(); ++i) {
+    weights.data[i] = half(static_cast<float>(i));
+  }
+
+  for (auto storage : env->GetSupportedStorages()) {
+    for (auto precision : env->GetSupportedPrecisions()) {
+      for (auto weights_layout :
+           {WeightsLayout::kOHWIOGroupI4O4, WeightsLayout::kOHWIOGroupO4I4}) {
+        conv_weight_desc.layout = weights_layout;
+        OperationDef op_def;
+        op_def.precision = precision;
+        auto data_type = DeduceDataTypeFromPrecision(precision);
+        op_def.src_tensors.push_back({data_type, storage, Layout::BHWC});
+        op_def.dst_tensors.push_back(
+            {data_type, TensorStorageType::BUFFER, Layout::UNKNOWN});
+        RETURN_IF_ERROR(ConvolutionWeightsConverterTest(
+            weights, conv_weight_desc, env, op_def));
+      }
+    }
+  }
+  return absl::OkStatus();
+}
+
+absl::Status ConverterToConvTransposedWeights4x4Test(
+    TestExecutionEnvironment* env) {
+  const int kSrcChannels = 7;
+  const int kDstChannels = 11;
+  auto weights_shape = OHWI(kDstChannels, 4, 4, kSrcChannels);
+  WeightsDescription weight_desc;
+  weight_desc.spatial_remap = {10, 11, 14, 15, 8, 9, 12, 13,
+                               2,  3,  6,  7,  0, 1, 4,  5};
+
+  Tensor<OHWI, DataType::FLOAT32> weights;
+  weights.shape = weights_shape;
+  weights.data.resize(weights_shape.DimensionsProduct());
+  for (int i = 0; i < weights.data.size(); ++i) {
+    weights.data[i] = half(static_cast<float>(i));
+  }
+
+  for (auto storage : env->GetSupportedStorages()) {
+    for (auto precision : env->GetSupportedPrecisions()) {
+      for (auto weights_layout : {WeightsLayout::kOICustomSpatialI4O4,
+                                  WeightsLayout::kOICustomSpatialO4I4}) {
+        weight_desc.layout = weights_layout;
+        OperationDef op_def;
+        op_def.precision = precision;
+        auto data_type = DeduceDataTypeFromPrecision(precision);
+        op_def.src_tensors.push_back({data_type, storage, Layout::BHWC});
+        op_def.dst_tensors.push_back(
+            {data_type, TensorStorageType::BUFFER, Layout::UNKNOWN});
+        RETURN_IF_ERROR(
+            ConvolutionWeightsConverterTest(weights, weight_desc, env, op_def));
+      }
+    }
+  }
+  return absl::OkStatus();
+}
+
+}  // namespace gpu
+}  // namespace tflite

tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter_test_util.h

@@ -0,0 +1,36 @@
+/* Copyright 2021 The TensorFlow Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+==============================================================================*/
+
+#ifndef TENSORFLOW_LITE_DELEGATES_GPU_COMMON_TASKS_CONV_WEIGHTS_CONVERTER_TEST_UTIL_H_
+#define TENSORFLOW_LITE_DELEGATES_GPU_COMMON_TASKS_CONV_WEIGHTS_CONVERTER_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 ConverterToConvWeights1x1OutX4Test(TestExecutionEnvironment* env);
+absl::Status ConverterToConvWeights1x1OutX4UnalignedTest(
+    TestExecutionEnvironment* env);
+absl::Status ConverterToConvWeights1x1OutX2Test(TestExecutionEnvironment* env);
+absl::Status ConverterToConvWeightsOutX2Test(TestExecutionEnvironment* env);
+absl::Status ConverterToConvTransposedWeights4x4Test(
+    TestExecutionEnvironment* env);
+
+}  // namespace gpu
+}  // namespace tflite
+
+#endif  // TENSORFLOW_LITE_DELEGATES_GPU_COMMON_TASKS_CONV_WEIGHTS_CONVERTER_TEST_UTIL_H_

tensorflow/lite/delegates/gpu/metal/kernels/BUILD

@@ -112,6 +112,29 @@ ios_unit_test(
     deps = [":conv_test_lib"],
 )
 
+objc_library(
+    name = "conv_weights_converter_test_lib",
+    testonly = 1,
+    srcs = ["conv_weights_converter_test.mm"],
+    sdk_frameworks = ["XCTest"],
+    deps = [
+        ":test_util",
+        "//tensorflow/lite/delegates/gpu/common/tasks:conv_weights_converter_test_util",
+    ],
+)
+
+ios_unit_test(
+    name = "conv_weights_converter_test",
+    testonly = 1,
+    minimum_os_version = "11.0",
+    runner = tflite_ios_lab_runner("IOS_LATEST"),
+    tags = tf_gpu_tests_tags() + [
+        "notap",
+        "tflite_not_portable_android",
+    ],
+    deps = [":conv_weights_converter_test_lib"],
+)
+
 objc_library(
     name = "depthwise_conv_test_lib",
     testonly = 1,
@@ -666,6 +689,7 @@ objc_library(
         "add_test.mm",
         "concat_test.mm",
         "conv_test.mm",
+        "conv_weights_converter_test.mm",
         "depthwise_conv_test.mm",
         "elementwise_test.mm",
         "fully_connected_test.mm",
@@ -702,6 +726,7 @@ objc_library(
         "//tensorflow/lite/delegates/gpu/common:util",
         "//tensorflow/lite/delegates/gpu/common/tasks:add_test_util",
         "//tensorflow/lite/delegates/gpu/common/tasks:concat_test_util",
+        "//tensorflow/lite/delegates/gpu/common/tasks:conv_weights_converter_test_util",
         "//tensorflow/lite/delegates/gpu/common/tasks:convolution_transposed_4x4_test_util",
         "//tensorflow/lite/delegates/gpu/common/tasks:convolution_transposed_test_util",
         "//tensorflow/lite/delegates/gpu/common/tasks:depthwise_conv_3x3_stride_h2_test_util",

tensorflow/lite/delegates/gpu/metal/kernels/conv_weights_converter_test.mm

@@ -0,0 +1,54 @@
+/* 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.
+==============================================================================*/
+
+#import <XCTest/XCTest.h>
+
+#include "tensorflow/lite/delegates/gpu/common/status.h"
+#include "tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter_test_util.h"
+#include "tensorflow/lite/delegates/gpu/metal/kernels/test_util.h"
+
+@interface ConvWeightsConverterMetalTest : XCTestCase
+@end
+
+@implementation ConvWeightsConverterMetalTest {
+  tflite::gpu::metal::MetalExecutionEnvironment exec_env_;
+}
+
+- (void)testConverterToConvWeights1x1OutX4 {
+  const auto status = ConverterToConvWeights1x1OutX4Test(&exec_env_);
+  XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
+}
+
+- (void)testConverterToConvWeights1x1OutX4Unaligned {
+  const auto status = ConverterToConvWeights1x1OutX4UnalignedTest(&exec_env_);
+  XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
+}
+
+- (void)testConverterToConvWeights1x1OutX2 {
+  const auto status = ConverterToConvWeights1x1OutX2Test(&exec_env_);
+  XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
+}
+
+- (void)testConverterToConvWeightsOutX2 {
+  const auto status = ConverterToConvWeightsOutX2Test(&exec_env_);
+  XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
+}
+
+- (void)testConverterToConvTransposedWeights4x4 {
+  const auto status = ConverterToConvTransposedWeights4x4Test(&exec_env_);
+  XCTAssertTrue(status.ok(), @"%s", std::string(status.message()).c_str());
+}
+
+@end