From e1b78c1f496353bf0520ac7f86222da320092559 Mon Sep 17 00:00:00 2001 From: chenzupeng Date: Wed, 12 Aug 2020 19:41:39 +0800 Subject: [PATCH] add opencl reshape op --- .../src/runtime/kernel/opencl/CMakeLists.txt | 1 + .../runtime/kernel/opencl/cl/fp16/reshape.cl | 14 +++ .../runtime/kernel/opencl/cl/fp32/reshape.cl | 14 +++ .../kernel/opencl/kernel/convolution.cc | 3 +- .../runtime/kernel/opencl/kernel/reshape.cc | 114 ++++++++++++++++++ .../runtime/kernel/opencl/kernel/reshape.h | 44 +++++++ mindspore/lite/test/CMakeLists.txt | 1 + 7 files changed, 190 insertions(+), 1 deletion(-) create mode 100644 mindspore/lite/src/runtime/kernel/opencl/cl/fp16/reshape.cl create mode 100644 mindspore/lite/src/runtime/kernel/opencl/cl/fp32/reshape.cl create mode 100644 mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc create mode 100644 mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.h diff --git a/mindspore/lite/src/runtime/kernel/opencl/CMakeLists.txt b/mindspore/lite/src/runtime/kernel/opencl/CMakeLists.txt index b090065ca..023973188 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/CMakeLists.txt +++ b/mindspore/lite/src/runtime/kernel/opencl/CMakeLists.txt @@ -10,4 +10,5 @@ set(OPENCL_KERNEL_SRC ${CMAKE_CURRENT_SOURCE_DIR}/kernel/concat.cc ${CMAKE_CURRENT_SOURCE_DIR}/kernel/conv2d_transpose.cc ${CMAKE_CURRENT_SOURCE_DIR}/kernel/transpose.cc + ${CMAKE_CURRENT_SOURCE_DIR}/kernel/reshape.cc ) diff --git a/mindspore/lite/src/runtime/kernel/opencl/cl/fp16/reshape.cl b/mindspore/lite/src/runtime/kernel/opencl/cl/fp16/reshape.cl new file mode 100644 index 000000000..867d1e3d1 --- /dev/null +++ b/mindspore/lite/src/runtime/kernel/opencl/cl/fp16/reshape.cl @@ -0,0 +1,14 @@ +#define FLT half +#define FLT4 half4 +#define READ_IMAGE read_imageh +#define WRITE_IMAGE write_imageh +__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; +__kernel void reshape(__read_only image2d_t src_data, __write_only image2d_t dst_data, int4 size) { + int X = get_global_id(0); + int Y = get_global_id(1); + int Z = get_global_id(2); + if (X >= size.x || Y >= size.y || Z >= size.z) { + return; + } + WRITE_IMAGE(dst_data, (int2)(Y * size.z + Z, X), READ_IMAGE(src_data, smp_zero, (int2)(Y * size.z + Z, X))); +} diff --git a/mindspore/lite/src/runtime/kernel/opencl/cl/fp32/reshape.cl b/mindspore/lite/src/runtime/kernel/opencl/cl/fp32/reshape.cl new file mode 100644 index 000000000..e752c4b3f --- /dev/null +++ b/mindspore/lite/src/runtime/kernel/opencl/cl/fp32/reshape.cl @@ -0,0 +1,14 @@ +#define FLT float +#define FLT4 float4 +#define READ_IMAGE read_imagef +#define WRITE_IMAGE write_imagef +__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; +__kernel void reshape(__read_only image2d_t src_data, __write_only image2d_t dst_data, int4 size) { + int X = get_global_id(0); + int Y = get_global_id(1); + int Z = get_global_id(2); + if (X >= size.x || Y >= size.y || Z >= size.z) { + return; + } + WRITE_IMAGE(dst_data, (int2)(Y * size.z + Z, X), READ_IMAGE(src_data, smp_zero, (int2)(Y * size.z + Z, X))); +} diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc index 91ae6808e..3d88a353b 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc @@ -30,7 +30,6 @@ namespace mindspore::kernel { int ConvolutionOpenCLKernel::Init() { static int count = 0; - std::cout << "ConvolutionOpenCLKernel::Init()\n"; std::set build_options; std::string source = CodeGen(); std::string program_name = "convolution" + std::to_string(count); @@ -41,6 +40,8 @@ int ConvolutionOpenCLKernel::Init() { ocl_runtime->LoadSource(program_name, source); ocl_runtime->BuildKernel(kernel_, program_name, kernel_name, build_options); this->InitBuffer(); + out_tensors_[0]->SetFormat(schema::Format_NHWC4); + MS_LOG(DEBUG) << kernel_name << " Init Done!"; return RET_OK; } diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc new file mode 100644 index 000000000..4394cde16 --- /dev/null +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc @@ -0,0 +1,114 @@ +/** + * Copyright 2019 Huawei Technologies Co., Ltd + * + * 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 +#include +#include "include/errorcode.h" +#include "src/kernel_registry.h" +#include "src/runtime/opencl/opencl_runtime.h" +#include "src/runtime/kernel/opencl/kernel/reshape.h" +#include "src/runtime/kernel/opencl/cl/fp16/reshape.cl.inc" +#include "src/runtime/kernel/opencl/cl/fp32/reshape.cl.inc" + +using mindspore::kernel::KERNEL_ARCH::kGPU; +using mindspore::lite::KernelRegistrar; +using mindspore::lite::RET_ERROR; +using mindspore::lite::RET_OK; +using mindspore::schema::PrimitiveType_Reshape; + +namespace mindspore::kernel { + +int ReshapeOpenCLKernel::Init() { + std::string kernel_name = "reshape"; + auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance(); + +#ifdef PROGRAM_WITH_IL + ocl_runtime->CreateKernelFromIL(kernel_(), kernel_name); +#else + std::set build_options; +#ifdef ENABLE_FP16 + std::string source = reshape_source_fp16; +#else + std::string source = reshape_source_fp32; +#endif + std::string program_name = "reshape"; + ocl_runtime->LoadSource(program_name, source); + ocl_runtime->BuildKernel(kernel_, program_name, kernel_name, build_options); +#endif + out_tensors_[0]->SetFormat(schema::Format_NHWC); + MS_LOG(DEBUG) << kernel_name << " Init Done!"; + return RET_OK; +} + +int ReshapeOpenCLKernel::ReSize() { return 0; } + +int ReshapeOpenCLKernel::GetImageSize(size_t idx, std::vector *img_size) { + size_t im_dst_x, im_dst_y; + std::vector shapex = in_tensors_[0]->shape(); + int h = shapex[1]; + int w = shapex[2]; + int c = shapex[3]; + im_dst_x = UP_DIV(w * c, C4NUM); + im_dst_y = h; +#ifdef ENABLE_FP16 + size_t img_dtype = CL_HALF_FLOAT; +#else + size_t img_dtype = CL_FLOAT; +#endif + img_size->clear(); + std::vector vec{im_dst_x, im_dst_y, img_dtype}; + *img_size = vec; + return RET_OK; +} + +int ReshapeOpenCLKernel::Run() { + MS_LOG(DEBUG) << this->name() << " Running!"; + std::vector shapex = in_tensors_[0]->shape(); + int h = shapex[1]; + int w = shapex[2]; + int c = shapex[3]; + int c4 = UP_DIV(c, C4NUM); + auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance(); + // local size should less than MAX_GROUP_SIZE + std::vector local = {}; + std::vector global = {(size_t)h, (size_t)w, (size_t)c4}; + cl_int4 size = {h, w, c4, 1}; + ocl_runtime->SetKernelArg(kernel_, 0, in_tensors_[0]->Data()); + ocl_runtime->SetKernelArg(kernel_, 1, out_tensors_[0]->Data()); + ocl_runtime->SetKernelArg(kernel_, 2, size); + ocl_runtime->RunKernel(kernel_, global, local, nullptr); + return RET_OK; +} + +kernel::LiteKernel *OpenCLReshapeKernelCreator(const std::vector &inputs, + const std::vector &outputs, + OpParameter *opParameter, const lite::Context *ctx, + const kernel::KernelKey &desc, const lite::Primitive *primitive) { + auto *kernel = new (std::nothrow) ReshapeOpenCLKernel(reinterpret_cast(opParameter), inputs, outputs); + if (kernel == nullptr) { + MS_LOG(ERROR) << "kernel " << opParameter->name_ << " create failed."; + return nullptr; + } + auto ret = kernel->Init(); + if (ret != RET_OK) { + delete kernel; + return nullptr; + } + return kernel; +} + +REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_Reshape, OpenCLReshapeKernelCreator) +} // namespace mindspore::kernel diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.h b/mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.h new file mode 100644 index 000000000..4a35c4488 --- /dev/null +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.h @@ -0,0 +1,44 @@ +/** + * Copyright 2020 Huawei Technologies Co., Ltd + * + * 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 MINDSPORE_LITE_SRC_BACKEND_OPENCL_RESHAPE_H_ +#define MINDSPORE_LITE_SRC_BACKEND_OPENCL_RESHAPE_H_ + +#include + +#include "src/lite_kernel.h" +#include "src/runtime/opencl/opencl_runtime.h" +#include "src/runtime/kernel/opencl/opencl_kernel.h" + +namespace mindspore::kernel { +class ReshapeOpenCLKernel : public OpenCLKernel { + public: + explicit ReshapeOpenCLKernel(OpParameter *parameter, const std::vector &inputs, + const std::vector &outputs) + : OpenCLKernel(parameter, inputs, outputs) {} + ~ReshapeOpenCLKernel() override{}; + + int Init() override; + int ReSize() override; + int Run() override; + int GetImageSize(size_t idx, std::vector *img_size) override; + + private: + cl::Kernel kernel_; +}; +} // namespace mindspore::kernel + +#endif // MINDSPORE_LITE_SRC_BACKEND_OPENCL_RESHAPE_H_ diff --git a/mindspore/lite/test/CMakeLists.txt b/mindspore/lite/test/CMakeLists.txt index 6c139fc9b..08e0930c3 100644 --- a/mindspore/lite/test/CMakeLists.txt +++ b/mindspore/lite/test/CMakeLists.txt @@ -145,6 +145,7 @@ if (SUPPORT_GPU) # ${LITE_DIR}/src/runtime/kernel/opencl/kernel/leaky_relu.cc ${LITE_DIR}/src/runtime/kernel/opencl/kernel/conv2d_transpose.cc ${LITE_DIR}/src/runtime/kernel/opencl/kernel/transpose.cc + ${LITE_DIR}/src/runtime/kernel/opencl/kernel/reshape.cc ) endif() ### minddata lite -- GitLab