diff --git a/tensorflow/lite/delegates/gpu/cl/kernels/BUILD b/tensorflow/lite/delegates/gpu/cl/kernels/BUILD
index 88d837998c4..15fb046b8cd 100644
--- a/tensorflow/lite/delegates/gpu/cl/kernels/BUILD
+++ b/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",
diff --git a/tensorflow/lite/delegates/gpu/cl/kernels/conv_weights_converter_test.cc b/tensorflow/lite/delegates/gpu/cl/kernels/conv_weights_converter_test.cc
new file mode 100644
index 00000000000..34789b9a695
--- /dev/null
+++ b/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
diff --git a/tensorflow/lite/delegates/gpu/common/tasks/BUILD b/tensorflow/lite/delegates/gpu/common/tasks/BUILD
index 862915f2461..c086937d557 100644
--- a/tensorflow/lite/delegates/gpu/common/tasks/BUILD
+++ b/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"],
diff --git a/tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter_test_util.cc b/tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter_test_util.cc
new file mode 100644
index 00000000000..790cabd6bc3
--- /dev/null
+++ b/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
diff --git a/tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter_test_util.h b/tensorflow/lite/delegates/gpu/common/tasks/conv_weights_converter_test_util.h
new file mode 100644
index 00000000000..1afa80e9b7b
--- /dev/null
+++ b/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_
diff --git a/tensorflow/lite/delegates/gpu/metal/kernels/BUILD b/tensorflow/lite/delegates/gpu/metal/kernels/BUILD
index 2e9524289bc..7360dedebe4 100644
--- a/tensorflow/lite/delegates/gpu/metal/kernels/BUILD
+++ b/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",
diff --git a/tensorflow/lite/delegates/gpu/metal/kernels/conv_weights_converter_test.mm b/tensorflow/lite/delegates/gpu/metal/kernels/conv_weights_converter_test.mm
new file mode 100644
index 00000000000..2abcfde99c0
--- /dev/null
+++ b/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