From 0e8b30627b32add41713cff3a4eb8cf62a6c0855 Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Wed, 21 Aug 2019 18:10:29 -0700 Subject: [PATCH] Implement OpenGL converters using new SPI. OpenGL API2 implementation will use new converters. PiperOrigin-RevId: 264736807 --- .../lite/delegates/gpu/gl/kernels/BUILD | 46 ++ .../delegates/gpu/gl/kernels/converter.cc | 395 ++++++++++++++++++ .../lite/delegates/gpu/gl/kernels/converter.h | 37 ++ .../gpu/gl/kernels/converter_test.cc | 166 ++++++++ tensorflow/lite/delegates/gpu/spi.h | 12 + 5 files changed, 656 insertions(+) create mode 100644 tensorflow/lite/delegates/gpu/gl/kernels/converter.cc create mode 100644 tensorflow/lite/delegates/gpu/gl/kernels/converter.h create mode 100644 tensorflow/lite/delegates/gpu/gl/kernels/converter_test.cc diff --git a/tensorflow/lite/delegates/gpu/gl/kernels/BUILD b/tensorflow/lite/delegates/gpu/gl/kernels/BUILD index 63b068312ec..03f8a479964 100644 --- a/tensorflow/lite/delegates/gpu/gl/kernels/BUILD +++ b/tensorflow/lite/delegates/gpu/gl/kernels/BUILD @@ -5,6 +5,52 @@ package( licenses = ["notice"], # Apache 2.0 ) +cc_library( + name = "converter", + srcs = ["converter.cc"], + hdrs = ["converter.h"], + deps = [ + "//tensorflow/lite/delegates/gpu:spi", + "//tensorflow/lite/delegates/gpu/common:shape", + "//tensorflow/lite/delegates/gpu/common:status", + "//tensorflow/lite/delegates/gpu/common:types", + "//tensorflow/lite/delegates/gpu/common:util", + "//tensorflow/lite/delegates/gpu/gl:command_queue", + "//tensorflow/lite/delegates/gpu/gl:gl_buffer", + "//tensorflow/lite/delegates/gpu/gl:gl_program", + "//tensorflow/lite/delegates/gpu/gl:gl_shader", + "@com_google_absl//absl/strings", + "@com_google_absl//absl/types:span", + ], +) + +cc_test( + name = "converter_test", + size = "small", + srcs = ["converter_test.cc"], + linkopts = [ + "-lEGL", + "-lGLESv3", + ], + tags = [ + "local", + "nobuilder", + "notap", + "tflite_not_portable_ios", + ], + deps = [ + ":converter", + "//tensorflow/lite/delegates/gpu/common:convert", + "//tensorflow/lite/delegates/gpu/common:shape", + "//tensorflow/lite/delegates/gpu/common:status", + "//tensorflow/lite/delegates/gpu/gl:egl_environment", + "//tensorflow/lite/delegates/gpu/gl:gl_buffer", + "//tensorflow/lite/delegates/gpu/gl:portable", + "@com_google_absl//absl/types:span", + "@com_google_googletest//:gtest_main", + ], +) + cc_library( name = "add", srcs = ["add.cc"], diff --git a/tensorflow/lite/delegates/gpu/gl/kernels/converter.cc b/tensorflow/lite/delegates/gpu/gl/kernels/converter.cc new file mode 100644 index 00000000000..a919b18402e --- /dev/null +++ b/tensorflow/lite/delegates/gpu/gl/kernels/converter.cc @@ -0,0 +1,395 @@ +/* Copyright 2019 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/gl/kernels/converter.h" + +#include "absl/strings/str_cat.h" +#include "absl/types/span.h" +#include "tensorflow/lite/delegates/gpu/common/shape.h" +#include "tensorflow/lite/delegates/gpu/common/status.h" +#include "tensorflow/lite/delegates/gpu/common/types.h" +#include "tensorflow/lite/delegates/gpu/common/util.h" +#include "tensorflow/lite/delegates/gpu/gl/gl_buffer.h" +#include "tensorflow/lite/delegates/gpu/gl/gl_program.h" +#include "tensorflow/lite/delegates/gpu/gl/gl_shader.h" + +namespace tflite { +namespace gpu { +namespace gl { +namespace { + +// Wraps given SSBO into GlBuffer object that does not have ownership. +Status WrapSSBO(OpenGlBuffer ssbo, GlBuffer* buffer) { + int64_t size_bytes; + RETURN_IF_ERROR(GetSSBOSize(ssbo.id, &size_bytes)); + *buffer = GlBuffer(GL_SHADER_STORAGE_BUFFER, ssbo.id, size_bytes, 0, false); + return OkStatus(); +} + +std::string GetShaderHeader(const uint3& localsize) { + return absl::StrCat("#version 310 es\nlayout(local_size_x = ", localsize.x, + ", local_size_y = ", localsize.y, + ", local_size_z = ", localsize.z, ") in;\n"); +} + +class OpenGlConverterImpl : public TensorObjectConverter { + public: + explicit OpenGlConverterImpl(CommandQueue* command_queue) + : command_queue_(command_queue) {} + + virtual Status Init(const TensorObjectDef& input_def, + const TensorObjectDef& output_def) = 0; + + protected: + Status InitializeProgram(const uint3& workgroup_size, + const std::string& shader_source) { + workgroup_size_ = workgroup_size; + GlShader shader; + RETURN_IF_ERROR(GlShader::CompileShader( + GL_COMPUTE_SHADER, GetShaderHeader(workgroup_size) + shader_source, + &shader)); + return GlProgram::CreateWithShader(shader, &program_); + } + + Status Dispatch(const uint3& workload) { + uint3 num_workgroups = IntegralDivideRoundUp(workload, workgroup_size_); + if (command_queue_) { + return command_queue_->Dispatch(program_, num_workgroups); + } + return program_.Dispatch(num_workgroups); + } + + GlProgram program_; + uint3 workgroup_size_; + CommandQueue* command_queue_; +}; + +bool IsSupportedDataType(DataType type) { return type == DataType::FLOAT32; } + +uint32_t SizeInBytesDHWC4(const BHWC& shape) { + return shape.b * shape.h * shape.w * AlignByN(shape.c, 4) * sizeof(float); +} + +uint32_t SizeInBytesBHWC(const BHWC& shape) { + return shape.DimensionsProduct() * sizeof(float); +} + +// Implements conversion from OpenGL-specific tensor layout to BHWC. +class FromTensorConverter : public OpenGlConverterImpl { + public: + explicit FromTensorConverter(CommandQueue* command_queue) + : OpenGlConverterImpl(command_queue) {} + + static bool IsSupported(const ObjectDef& input, const ObjectDef& output) { + return IsSupportedDataType(input.data_type) && + IsSupportedDataType(output.data_type) && + // Output is always SSBO/BHWC + output.object_type == ObjectType::OPENGL_SSBO && + output.data_layout == DataLayout::BHWC && + // SSBO/DHWC4 -> + input.object_type == ObjectType::OPENGL_SSBO && + input.data_layout == DataLayout::DHWC4; + } + + Status Init(const TensorObjectDef& input_def, + const TensorObjectDef& output_def) final { + shape_ = BHWC(output_def.dimensions.b, output_def.dimensions.h, + output_def.dimensions.w, output_def.dimensions.c); + if (shape_.b != 1) { + return UnimplementedError( + "FromTensorConverter: Batch size != 1 is not supported."); + } + + return InitializeProgram(uint3(8, 4, 2), R"( + layout(std430) buffer; + precision highp float; + + layout(binding = 0) readonly buffer B0 { + vec4 elements[]; + } input_data; + + layout(binding = 1) writeonly buffer B1 { + float elements[]; + } output_data; + + uniform ivec4 sizes; + + void main() { + ivec3 gid = ivec3(gl_GlobalInvocationID.xyz); + if (gid.x >= sizes.x || gid.y >= sizes.y || gid.z >= sizes.z) { + return; + } + output_data.elements[(gid.y * sizes.x + gid.x) * sizes.z + gid.z] = input_data.elements[(gid.z / 4 * sizes.y + gid.y) * sizes.x + gid.x][gid.z % 4]; + })"); + } + + Status Convert(const TensorObject& input_obj, + const TensorObject& output_obj) override { + auto output = absl::get_if(&output_obj); + if (!output || !output->id) { + return InvalidArgumentError("Missing output in converter"); + } + auto input = absl::get_if(&input_obj); + if (!input || !input->id) { + return InvalidArgumentError("Missing input in converter"); + } + if (input->id == output->id) { + return InvalidArgumentError("Can not execute inplace conversion"); + } + GlBuffer input_ssbo; + RETURN_IF_ERROR(WrapSSBO(*input, &input_ssbo)); + GlBuffer output_ssbo; + RETURN_IF_ERROR(WrapSSBO(*output, &output_ssbo)); + + if (input_ssbo.bytes_size() != SizeInBytesDHWC4(shape_)) { + return InvalidArgumentError( + "FromTensorConverter: input data size does not match expected size."); + } + if (output_ssbo.bytes_size() != SizeInBytesBHWC(shape_)) { + return InvalidArgumentError( + "FromTensorConverter: output data size does not match expected " + "size."); + } + RETURN_IF_ERROR(program_.SetParameter( + {"sizes", + int4(static_cast(shape_.w), static_cast(shape_.h), + static_cast(shape_.c), 0)})); + RETURN_IF_ERROR(input_ssbo.BindToIndex(0)); + RETURN_IF_ERROR(output_ssbo.BindToIndex(1)); + return Dispatch(uint3(shape_.w, shape_.h, shape_.c)); + } + + BHWC shape_; +}; + +// Implements conversion from BHWC to OpenCL-specific tensor layout. +class ToTensorConverter : public OpenGlConverterImpl { + public: + explicit ToTensorConverter(CommandQueue* command_queue) + : OpenGlConverterImpl(command_queue) {} + + static bool IsSupported(const ObjectDef& input, const ObjectDef& output) { + return IsSupportedDataType(input.data_type) && + IsSupportedDataType(output.data_type) && + // Input is always SSBO/BHWC + input.object_type == ObjectType::OPENGL_SSBO && + input.data_layout == DataLayout::BHWC && + // -> SSBO/DHWC4 + output.object_type == ObjectType::OPENGL_SSBO && + output.data_layout == DataLayout::DHWC4; + } + + Status Init(const TensorObjectDef& input_def, + const TensorObjectDef& output_def) final { + shape_ = BHWC(output_def.dimensions.b, output_def.dimensions.h, + output_def.dimensions.w, output_def.dimensions.c); + if (shape_.b != 1) { + return UnimplementedError( + "FromTensorConverter: Batch size != 1 is not supported."); + } + + return InitializeProgram(uint3(8, 4, 2), R"( + layout(std430) buffer; + precision highp float; + + layout(binding = 0) readonly buffer B0 { + float elements[]; + } input_data; + + layout(binding = 1) writeonly buffer B1 { + vec4 elements[]; + } output_data; + + uniform ivec4 sizes; + + void main() { + ivec3 gid = ivec3(gl_GlobalInvocationID.xyz); + if (gid.x >= sizes.x || gid.y >= sizes.y || gid.z >= sizes.w) { + return; + } + vec4 v = vec4(0); + int dst_channel = gid.z * 4; + int index = (gid.y * sizes.x + gid.x) * sizes.z + dst_channel; + for (int i = 0; i < 4; ++i, ++index, ++dst_channel) { + if (dst_channel >= sizes.z) break; + v[i] = input_data.elements[index]; + } + output_data.elements[(gid.z * sizes.y + gid.y) * sizes.x + gid.x] = v; + })"); + } + + Status Convert(const TensorObject& input_obj, + const TensorObject& output_obj) override { + auto output = absl::get_if(&output_obj); + if (!output || !output->id) { + return InvalidArgumentError("Missing output in converter"); + } + auto input = absl::get_if(&input_obj); + if (!input || !input->id) { + return InvalidArgumentError("Missing input in converter"); + } + if (input->id == output->id) { + return InvalidArgumentError("Can not execute inplace conversion"); + } + GlBuffer input_ssbo; + RETURN_IF_ERROR(WrapSSBO(*input, &input_ssbo)); + GlBuffer output_ssbo; + RETURN_IF_ERROR(WrapSSBO(*output, &output_ssbo)); + + if (input_ssbo.bytes_size() != SizeInBytesBHWC(shape_)) { + return InvalidArgumentError( + "ToTensorConverter: input data size does not match expected size."); + } + if (output_ssbo.bytes_size() != SizeInBytesDHWC4(shape_)) { + return InvalidArgumentError( + "ToTensorConverter: output data size does not match expected size."); + } + auto d = IntegralDivideRoundUp(shape_.c, 4); + RETURN_IF_ERROR(program_.SetParameter( + {"sizes", + int4(static_cast(shape_.w), static_cast(shape_.h), + static_cast(shape_.c), static_cast(d))})); + RETURN_IF_ERROR(input_ssbo.BindToIndex(0)); + RETURN_IF_ERROR(output_ssbo.BindToIndex(1)); + return Dispatch(uint3(shape_.w, shape_.h, d)); + } + + BHWC shape_; +}; + +// Copies data from one object of the same type and layout to another object. +class TrivialCopier : public TensorObjectConverter { + public: + static bool IsSupported(const ObjectDef& input, const ObjectDef& output) { + return input.object_type == ObjectType::OPENGL_SSBO && + input.data_type == output.data_type && + input.object_type == output.object_type && + input.data_layout == output.data_layout; + } + + Status Convert(const TensorObject& input_obj, + const TensorObject& output_obj) override { + auto ssbo_input = absl::get_if(&input_obj); + auto ssbo_output = absl::get_if(&output_obj); + if (ssbo_input && ssbo_output) { + return Copy(*ssbo_input, *ssbo_output); + } + return InternalError("Unexpected object"); + } + + Status Copy(OpenGlBuffer input, OpenGlBuffer output) { + if (input.id == output.id) { + return OkStatus(); + } + GlBuffer input_obj; + RETURN_IF_ERROR(WrapSSBO(input, &input_obj)); + GlBuffer output_obj; + RETURN_IF_ERROR(WrapSSBO(output, &output_obj)); + return CopyBuffer(input_obj, output_obj); + } +}; + +// Copies data from/to CPU into a tensor. +class CpuCopier : public TensorObjectConverter { + public: + static bool IsSupported(const ObjectDef& input, const ObjectDef& output) { + return input.data_type == output.data_type && + input.data_layout == output.data_layout && + ((input.object_type == ObjectType::CPU_MEMORY && + output.object_type == ObjectType::OPENGL_SSBO) || + (output.object_type == ObjectType::CPU_MEMORY && + input.object_type == ObjectType::OPENGL_SSBO)); + } + + Status Convert(const TensorObject& input_obj, + const TensorObject& output_obj) override { + auto cpu_input = absl::get_if(&input_obj); + auto cpu_output = absl::get_if(&output_obj); + if (cpu_input) { + auto ssbo_output = absl::get_if(&output_obj); + if (ssbo_output) { + GlBuffer gl_buffer; + RETURN_IF_ERROR(WrapSSBO(*ssbo_output, &gl_buffer)); + return gl_buffer.Write( + absl::MakeConstSpan(static_cast(cpu_input->data), + cpu_input->size_bytes)); + } + } else if (cpu_output) { + auto ssbo_input = absl::get_if(&input_obj); + if (ssbo_input) { + GlBuffer gl_buffer; + RETURN_IF_ERROR(WrapSSBO(*ssbo_input, &gl_buffer)); + return gl_buffer.Read(absl::MakeSpan( + static_cast(cpu_input->data), cpu_input->size_bytes)); + } + } + return InternalError("Unexpected object"); + } +}; + +class TensorConverterBuilderImpl : public TensorObjectConverterBuilder { + public: + explicit TensorConverterBuilderImpl(CommandQueue* command_queue) + : command_queue_(command_queue) {} + + bool IsSupported(const TensorObjectDef& input, + const TensorObjectDef& output) final { + const auto& input_def = input.object_def; + const auto& output_def = output.object_def; + return input.dimensions == output.dimensions && + (TrivialCopier::IsSupported(input_def, output_def) || + CpuCopier::IsSupported(input_def, output_def) || + FromTensorConverter::IsSupported(input_def, output_def) || + ToTensorConverter::IsSupported(input_def, output_def)); + } + + Status MakeConverter( + const TensorObjectDef& input, const TensorObjectDef& output, + std::unique_ptr* converter) final { + std::unique_ptr impl; + const auto& input_def = input.object_def; + const auto& output_def = output.object_def; + if (TrivialCopier::IsSupported(input_def, output_def)) { + *converter = absl::make_unique(); + return OkStatus(); + } else if (CpuCopier::IsSupported(input_def, output_def)) { + *converter = absl::make_unique(); + return OkStatus(); + } else if (FromTensorConverter::IsSupported(input_def, output_def)) { + impl = absl::make_unique(command_queue_); + } else if (ToTensorConverter::IsSupported(input_def, output_def)) { + impl = absl::make_unique(command_queue_); + } else { + return UnimplementedError("Unsupported conversion"); + } + RETURN_IF_ERROR(impl->Init(input, output)); + *converter = std::move(impl); + return OkStatus(); + } + + private: + CommandQueue* command_queue_; +}; + +} // namespace + +std::unique_ptr NewConverterBuilder( + CommandQueue* command_queue) { + return absl::make_unique(command_queue); +} + +} // namespace gl +} // namespace gpu +} // namespace tflite diff --git a/tensorflow/lite/delegates/gpu/gl/kernels/converter.h b/tensorflow/lite/delegates/gpu/gl/kernels/converter.h new file mode 100644 index 00000000000..c5f2ba208f8 --- /dev/null +++ b/tensorflow/lite/delegates/gpu/gl/kernels/converter.h @@ -0,0 +1,37 @@ +/* Copyright 2019 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_GL_KERNELS_CONVERTER_H_ +#define TENSORFLOW_LITE_DELEGATES_GPU_GL_KERNELS_CONVERTER_H_ + +#include + +#include "tensorflow/lite/delegates/gpu/gl/command_queue.h" +#include "tensorflow/lite/delegates/gpu/spi.h" + +namespace tflite { +namespace gpu { +namespace gl { + +// Supports conversions from DHWC4 to internal OpenGL tensor representation and +// back. Supports F32 only. +std::unique_ptr NewConverterBuilder( + CommandQueue* command_queue /* optional */); + +} // namespace gl +} // namespace gpu +} // namespace tflite + +#endif // TENSORFLOW_LITE_DELEGATES_GPU_GL_KERNELS_CONVERTER_H_ diff --git a/tensorflow/lite/delegates/gpu/gl/kernels/converter_test.cc b/tensorflow/lite/delegates/gpu/gl/kernels/converter_test.cc new file mode 100644 index 00000000000..daba2f6d9ef --- /dev/null +++ b/tensorflow/lite/delegates/gpu/gl/kernels/converter_test.cc @@ -0,0 +1,166 @@ +/* Copyright 2019 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/gl/kernels/converter.h" + +#include +#include + +#include +#include +#include "absl/types/span.h" +#include "tensorflow/lite/delegates/gpu/common/convert.h" +#include "tensorflow/lite/delegates/gpu/common/shape.h" +#include "tensorflow/lite/delegates/gpu/common/status.h" +#include "tensorflow/lite/delegates/gpu/gl/egl_environment.h" +#include "tensorflow/lite/delegates/gpu/gl/gl_buffer.h" +#include "tensorflow/lite/delegates/gpu/gl/portable_gl31.h" + +namespace tflite { +namespace gpu { +namespace gl { +namespace { + +inline std::vector GenerateFloats(float multiplier, int size) { + std::vector v(size); + for (int i = 0; i < size; ++i) { + v[i] = multiplier * i * (i % 2 == 0 ? -1 : 1); + } + return v; +} + +Dimensions ToDimensions(const BHWC& shape) { + return Dimensions(shape.b, shape.h, shape.w, shape.c); +} + +Status RunFromTensorTest(const BHWC& shape) { + // Create random input and calculate expected output for it. + std::vector input = + GenerateFloats(0.01, GetElementsSizeForPHWC4(shape)); + std::vector output(shape.DimensionsProduct(), 0); + RETURN_IF_ERROR( + ConvertFromPHWC4(absl::MakeConstSpan(input.data(), input.size()), shape, + absl::MakeSpan(output.data(), output.size()))); + + std::unique_ptr env; + RETURN_IF_ERROR(EglEnvironment::NewEglEnvironment(&env)); + + // Create input and output buffers + GlBuffer input_buffer; + RETURN_IF_ERROR(CreateReadOnlyShaderStorageBuffer( + absl::MakeConstSpan(input.data(), input.size()), &input_buffer)); + + GlBuffer output_buffer; + RETURN_IF_ERROR(CreateReadWriteShaderStorageBuffer( + shape.DimensionsProduct(), &output_buffer)); + + // Create converter and run it. + auto builder = NewConverterBuilder(nullptr); + TensorObjectDef input_def; + input_def.object_def.data_type = DataType::FLOAT32; + input_def.object_def.data_layout = DataLayout::DHWC4; + input_def.object_def.object_type = ObjectType::OPENGL_SSBO; + input_def.dimensions = ToDimensions(shape); + TensorObjectDef output_def = input_def; + output_def.object_def.data_layout = DataLayout::BHWC; + std::unique_ptr converter; + RETURN_IF_ERROR(builder->MakeConverter(input_def, output_def, &converter)); + RETURN_IF_ERROR(converter->Convert(OpenGlBuffer{input_buffer.id()}, + OpenGlBuffer{output_buffer.id()})); + + // Compare outputs. + std::vector converted_output(output.size(), 0); + RETURN_IF_ERROR(output_buffer.Read( + absl::MakeSpan(converted_output.data(), converted_output.size()))); + if (output != converted_output) { + return InternalError("Outputs don't match"); + } + return OkStatus(); +} + +TEST(FromTensor, Smoke) { + for (int32_t h : {1, 2, 3, 7, 20}) { + for (int32_t w : {1, 2, 4, 5, 11}) { + for (int32_t c : {1, 2, 4, 5, 8, 9}) { + BHWC shape(1, h, w, c); + auto status = RunFromTensorTest(shape); + EXPECT_TRUE(status.ok()) << status << ", shape = " << shape.h << " " + << shape.w << " " << shape.c; + } + } + } +} + +Status RunToTensorTest(const BHWC& shape) { + // Create random input and calculate expected output for it. + std::vector input = GenerateFloats(0.01, shape.DimensionsProduct()); + std::vector output(GetElementsSizeForPHWC4(shape), 0); + RETURN_IF_ERROR( + ConvertToPHWC4(absl::MakeConstSpan(input.data(), input.size()), shape, + absl::MakeSpan(output.data(), output.size()))); + + std::unique_ptr env; + RETURN_IF_ERROR(EglEnvironment::NewEglEnvironment(&env)); + + // Create input and output buffers + GlBuffer input_buffer; + RETURN_IF_ERROR(CreateReadOnlyShaderStorageBuffer( + absl::MakeConstSpan(input.data(), input.size()), &input_buffer)); + + GlBuffer output_buffer; + RETURN_IF_ERROR(CreateReadWriteShaderStorageBuffer( + GetElementsSizeForPHWC4(shape), &output_buffer)); + + // Create converter and run it. + auto builder = NewConverterBuilder(nullptr); + TensorObjectDef input_def; + input_def.object_def.data_type = DataType::FLOAT32; + input_def.object_def.data_layout = DataLayout::BHWC; + input_def.object_def.object_type = ObjectType::OPENGL_SSBO; + input_def.dimensions = ToDimensions(shape); + TensorObjectDef output_def = input_def; + output_def.object_def.data_layout = DataLayout::DHWC4; + std::unique_ptr converter; + RETURN_IF_ERROR(builder->MakeConverter(input_def, output_def, &converter)); + RETURN_IF_ERROR(converter->Convert(OpenGlBuffer{input_buffer.id()}, + OpenGlBuffer{output_buffer.id()})); + + // Compare outputs. + std::vector converted_output(output.size(), 0); + RETURN_IF_ERROR(output_buffer.Read( + absl::MakeSpan(converted_output.data(), converted_output.size()))); + if (output != converted_output) { + return InternalError("Outputs don't match"); + } + return OkStatus(); +} + +TEST(ToTensor, Smoke) { + for (int32_t h : {1, 2, 3, 7, 20}) { + for (int32_t w : {1, 2, 4, 5, 11}) { + for (int32_t c : {1, 2, 4, 5, 8, 9}) { + BHWC shape(1, h, w, c); + auto status = RunToTensorTest(shape); + EXPECT_TRUE(status.ok()) << status << ", shape = " << shape.h << " " + << shape.w << " " << shape.c; + } + } + } +} + +} // namespace +} // namespace gl +} // namespace gpu +} // namespace tflite diff --git a/tensorflow/lite/delegates/gpu/spi.h b/tensorflow/lite/delegates/gpu/spi.h index 023cc7a2c34..fcc3a5714ef 100644 --- a/tensorflow/lite/delegates/gpu/spi.h +++ b/tensorflow/lite/delegates/gpu/spi.h @@ -34,6 +34,18 @@ class TensorObjectConverter { const TensorObject& output) = 0; }; +class TensorObjectConverterBuilder { + public: + virtual ~TensorObjectConverterBuilder() = default; + + virtual bool IsSupported(const TensorObjectDef& input, + const TensorObjectDef& output) = 0; + + virtual Status MakeConverter( + const TensorObjectDef& input, const TensorObjectDef& output, + std::unique_ptr* converter) = 0; +}; + } // namespace gpu } // namespace tflite