!4357 add opencl reshape op

Merge pull request !4357 from chenzupeng/master-lite

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
     )

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)));
+}

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)));
+}

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<std::string> 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;
 }
 

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 <set>
+#include <string>
+#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<std::string> 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<size_t> *img_size) {
+  size_t im_dst_x, im_dst_y;
+  std::vector<int> 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<size_t> 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<int> 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<size_t> local = {};
+  std::vector<size_t> 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<lite::tensor::Tensor *> &inputs,
+                                               const std::vector<lite::tensor::Tensor *> &outputs,
+                                               OpParameter *opParameter, const lite::Context *ctx,
+                                               const kernel::KernelKey &desc, const lite::Primitive *primitive) {
+  auto *kernel = new (std::nothrow) ReshapeOpenCLKernel(reinterpret_cast<OpParameter *>(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

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 <vector>
+
+#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<lite::tensor::Tensor *> &inputs,
+                               const std::vector<lite::tensor::Tensor *> &outputs)
+      : OpenCLKernel(parameter, inputs, outputs) {}
+  ~ReshapeOpenCLKernel() override{};
+
+  int Init() override;
+  int ReSize() override;
+  int Run() override;
+  int GetImageSize(size_t idx, std::vector<size_t> *img_size) override;
+
+ private:
+  cl::Kernel kernel_;
+};
+}  // namespace mindspore::kernel
+
+#endif  // MINDSPORE_LITE_SRC_BACKEND_OPENCL_RESHAPE_H_

mindspore/lite/test/CMakeLists.txt

@@ -149,6 +149,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