Merge branch 'opencl' of https://github.com/PaddlePaddle/paddle-mobile into opencl

src/framework/cl/cl_scope.h

@@ -54,7 +54,7 @@ class CLScope {
     }
 
     auto program =
-        CLEngine::Instance()->CreateProgramWith(context_.get(), file_name);
+        CLEngine::Instance()->CreateProgramWith(context_.get(), "./cl_kernel/" + file_name);
     programs_[file_name] = std::move(program);
 
     status_ = clBuildProgram(program.get(), 0, 0, 0, 0, 0);

src/operators/kernel/cl/cl_kernel/common.h → src/operators/kernel/cl/cl_kernel/cl_common.h

@@ -14,7 +14,7 @@ limitations under the License. */
 
 #pragma once;
 
-/*
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 inline hafl4 activation(half4 in
 #ifdef PRELU
@@ -32,4 +32,3 @@ inline hafl4 activation(half4 in
   return output;
 }
 
-*/

src/operators/kernel/cl/cl_kernel/conv_kernel.inc.cl

@@ -23,9 +23,7 @@ conv_add_bn_relu
 
 */
 
-/*
-
-#include "common.h"
+#include "cl_common.h"
 
 __kernel void conv_1x1(__private const int global_size_dim0,
                        __private const int global_size_dim1,
@@ -43,8 +41,11 @@ __kernel void conv_1x1(__private const int global_size_dim0,
                        __private const int stride,
                        __private const int offset,
                        __private const int input_c,
+                       __private const int dilation,
                        __private const int input_width,/* of one block */
-                       __private const int input_height/* of one block */) {
+                       __private const int input_height,/* of one block */
+                       __private const int output_width,
+                       __private const int output_height) {
   const int out_c = get_global_id(0);
   const int out_w = get_global_id(1);
   const int out_nh = get_global_id(2);
@@ -114,7 +115,9 @@ __kernel void conv_3x3(__private const int global_size_dim0,
                                               __private const int input_c,
                                               __private const int dilation,
                                               __private const int input_width,/* of one block */
-                                              __private const int input_height/* of one block */) {
+                                              __private const int input_height,/* of one block */
+                                              __private const int output_width,
+                                              __private const int output_height) {
     int2 stride_xy = int2(stride, stride);
     int2 ouput_pos_in_one_block = int2(out_w, out_nh);
     int2 in_pos_in_one_block = ouput_pos_in_one_block * stride_xy + int2(offset, offset);
@@ -305,5 +308,3 @@ __kernel void depth_conv_3x3(__private const int global_size_dim0,
     int2 output_pos(out_c * global_size_dim1 + out_w, out_nh);
     write_imageh(output_image, output_pos, output);
 }
-
-*/

src/operators/kernel/cl/conv_add_bn_relu_kernel.cpp

@@ -109,8 +109,11 @@ void ConvAddBNReluKernel<GPU_CL, float>::Compute(
   int stride = param.Strides()[0];
   int offset = param.Offset();
   int input_c = param.Input()->CBlock();
+  int dilation = param.Dilations()[0];
   int input_width = param.Input()->WidthOfOneBlock();
   int input_height = param.Input()->HeightOfOneBlock();
+  int output_width = param.Output()->WidthOfOneBlock();
+  int output_height = param.Output()->HeightOfOneBlock();
 
   clSetKernelArg(kernel, 0, sizeof(int), &c_block);
   clSetKernelArg(kernel, 1, sizeof(int), &w);
@@ -124,8 +127,11 @@ void ConvAddBNReluKernel<GPU_CL, float>::Compute(
   clSetKernelArg(kernel, 9, sizeof(int), &stride);
   clSetKernelArg(kernel, 10, sizeof(int), &offset);
   clSetKernelArg(kernel, 11, sizeof(int), &input_c);
-  clSetKernelArg(kernel, 12, sizeof(int), &input_width);
-  clSetKernelArg(kernel, 13, sizeof(int), &input_height);
+  clSetKernelArg(kernel, 12, sizeof(int), &dilation);
+  clSetKernelArg(kernel, 13, sizeof(int), &input_width);
+  clSetKernelArg(kernel, 14, sizeof(int), &input_height);
+  clSetKernelArg(kernel, 15, sizeof(int), &output_width);
+  clSetKernelArg(kernel, 16, sizeof(int), &output_height);
 
   clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 3, NULL,
                          default_work_size.data(), NULL, 0, NULL, NULL);

src/operators/kernel/cl/conv_add_kernel.cpp

@@ -59,8 +59,11 @@ void ConvAddKernel<GPU_CL, float>::Compute(
   int stride = param.Strides()[0];
   int offset = param.Offset();
   int input_c = param.Input()->CBlock();
+  int dilation = param.Dilations()[0];
   int input_width = param.Input()->WidthOfOneBlock();
   int input_height = param.Input()->HeightOfOneBlock();
+  int output_width = param.Output()->WidthOfOneBlock();
+  int output_height = param.Output()->HeightOfOneBlock();
 
   clSetKernelArg(kernel, 0, sizeof(int), &c_block);
   clSetKernelArg(kernel, 1, sizeof(int), &w);
@@ -68,12 +71,15 @@ void ConvAddKernel<GPU_CL, float>::Compute(
   clSetKernelArg(kernel, 3, sizeof(cl_mem), &input);
   clSetKernelArg(kernel, 4, sizeof(cl_mem), &filter);
   clSetKernelArg(kernel, 5, sizeof(cl_mem), &biase);
-  clSetKernelArg(kernel, 8, sizeof(cl_mem), &output);
-  clSetKernelArg(kernel, 9, sizeof(int), &stride);
-  clSetKernelArg(kernel, 10, sizeof(int), &offset);
-  clSetKernelArg(kernel, 11, sizeof(int), &input_c);
-  clSetKernelArg(kernel, 12, sizeof(int), &input_width);
-  clSetKernelArg(kernel, 13, sizeof(int), &input_height);
+  clSetKernelArg(kernel, 6, sizeof(cl_mem), &output);
+  clSetKernelArg(kernel, 7, sizeof(int), &stride);
+  clSetKernelArg(kernel, 8, sizeof(int), &offset);
+  clSetKernelArg(kernel, 9, sizeof(int), &input_c);
+  clSetKernelArg(kernel, 10, sizeof(int), &dilation);
+  clSetKernelArg(kernel, 11, sizeof(int), &input_width);
+  clSetKernelArg(kernel, 12, sizeof(int), &input_height);
+  clSetKernelArg(kernel, 13, sizeof(int), &output_width);
+  clSetKernelArg(kernel, 14, sizeof(int), &output_height);
 
   clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 3, NULL,
                          default_work_size.data(), NULL, 0, NULL, NULL);

src/operators/kernel/cl/conv_kernel.cpp

@@ -60,6 +60,8 @@ void ConvKernel<GPU_CL, float>::Compute(const ConvParam<GPU_CL> &param) {
   int dilation = param.Dilations()[0];
   int input_width = param.Input()->WidthOfOneBlock();
   int input_height = param.Input()->HeightOfOneBlock();
+  int output_width = param.Output()->WidthOfOneBlock();
+  int output_height = param.Output()->HeightOfOneBlock();
 
   clSetKernelArg(kernel, 0, sizeof(int), &c_block);
   clSetKernelArg(kernel, 1, sizeof(int), &w);
@@ -73,14 +75,11 @@ void ConvKernel<GPU_CL, float>::Compute(const ConvParam<GPU_CL> &param) {
   clSetKernelArg(kernel, 9, sizeof(int), &dilation);
   clSetKernelArg(kernel, 10, sizeof(int), &input_width);
   clSetKernelArg(kernel, 11, sizeof(int), &input_height);
+  clSetKernelArg(kernel, 12, sizeof(int), &output_width);
+  clSetKernelArg(kernel, 13, sizeof(int), &output_height);
 
   clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 3, NULL,
                          default_work_size.data(), NULL, 0, NULL, NULL);
-
-  //  auto kernel = this->cl_helper_.KernelAt(0);
-  //  size_t global_work_size[3] = {1, 2, 3};
-  //  clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 3, NULL,
-  //  global_work_size, NULL, 0, NULL, NULL);
 }
 
 template class ConvKernel<GPU_CL, float>;

src/operators/kernel/cl/depthwise_conv_kernel.cpp

+/* Copyright (c) 2018 PaddlePaddle 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. */
+
+#ifdef DEPTHWISECONV_OP
+
+#include "operators/kernel/depthwise_conv_kernel.h"
+#include "operators/kernel/central-arm-func/depthwise_conv_arm_func.h"
+
+namespace paddle_mobile {
+namespace operators {
+
+template <>
+bool DepthwiseConvKernel<GPU_CL, float>::Init(ConvParam<GPU_CL> *param) {
+  DLOG << " depthwise conv kernel init begin ";
+  PADDLE_MOBILE_ENFORCE(
+          param->Filter()->dims()[2] == param->Filter()->dims()[3] &&
+          param->Paddings()[0] == param->Paddings()[1],
+          "need equal");
+  int offset = static_cast<int>(param->Filter()->dims()[2]) / 2 -
+               static_cast<int>(param->Paddings()[1]);
+  param->SetOffset(offset);
+  this->cl_helper_.AddKernel("depth_conv_3x3", "conv_add_bn_relu_kernel.cl");
+  DLOG << " depthwise conv kernel init end ";
+  return true;
+}
+
+template <>
+void DepthwiseConvKernel<GPU_CL, float>::Compute(const ConvParam<GPU_CL> &param) {
+  auto kernel = this->cl_helper_.KernelAt(0);
+  auto default_work_size = this->cl_helper_.DefaultWorkSize(*param.Output());
+  int c_block = default_work_size[0];
+  int w = default_work_size[1];
+  int nh = default_work_size[2];
+  auto input = param.Input()->GetCLImage();
+  auto filter = param.Filter()->GetCLImage();
+  auto output = param.Output();
+  int stride = param.Strides()[0];
+  int offset = param.Offset();
+  int input_c = param.Input()->CBlock();
+  int dilation = param.Dilations()[0];
+  int input_width = param.Input()->WidthOfOneBlock();
+  int input_height = param.Input()->HeightOfOneBlock();
+  int output_width = param.Output()->WidthOfOneBlock();
+  int output_height = param.Output()->HeightOfOneBlock();
+
+  clSetKernelArg(kernel, 0, sizeof(int), &c_block);
+  clSetKernelArg(kernel, 1, sizeof(int), &w);
+  clSetKernelArg(kernel, 2, sizeof(int), &nh);
+  clSetKernelArg(kernel, 3, sizeof(cl_mem), &input);
+  clSetKernelArg(kernel, 4, sizeof(cl_mem), &filter);
+  clSetKernelArg(kernel, 5, sizeof(cl_mem), &output);
+  clSetKernelArg(kernel, 6, sizeof(int), &stride);
+  clSetKernelArg(kernel, 7, sizeof(int), &offset);
+  clSetKernelArg(kernel, 8, sizeof(int), &input_c);
+  clSetKernelArg(kernel, 9, sizeof(int), &dilation);
+  clSetKernelArg(kernel, 10, sizeof(int), &input_width);
+  clSetKernelArg(kernel, 11, sizeof(int), &input_height);
+  clSetKernelArg(kernel, 12, sizeof(int), &output_width);
+  clSetKernelArg(kernel, 13, sizeof(int), &output_height);
+
+  clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 3, NULL,
+                         default_work_size.data(), NULL, 0, NULL, NULL);
+}
+
+template class DepthwiseConvKernel<GPU_CL, float>;
+
+}  // namespace operators
+}  // namespace paddle_mobile
+
+#endif
\ No newline at end of file

tools/build.sh

@@ -90,6 +90,8 @@ build_for_android() {
     fi
     cd "../build/release/${PLATFORM}"
     make -j 8
+    mkdir ./build/cl_kernel
+    cp ../../../src/operators/kernel/cl/cl_kernel/*  ./build/cl_kernel/
 }
 
 build_for_ios() {