diff --git a/src/common/types.h b/src/common/types.h index 6d38e4178907aa30968a6760a6ae5d69f4b61167..a5782e7394e78a6ccfe8d51da19b5da1caebdaed 100644 --- a/src/common/types.h +++ b/src/common/types.h @@ -39,7 +39,7 @@ struct PrecisionTrait { }; //! device type -enum DeviceTypeEnum { kINVALID = -1, kCPU = 0, kFPGA = 1, kGPU_MALI = 2 }; +enum DeviceTypeEnum { kINVALID = -1, kCPU = 0, kFPGA = 1, kGPU_MALI = 2, kGPU_CL = 3}; template struct DeviceType {}; @@ -47,6 +47,8 @@ struct DeviceType {}; typedef DeviceType CPU; typedef DeviceType FPGA; typedef DeviceType GPU_MALI; +typedef DeviceType GPU_CL; + //! data type enum DataType { diff --git a/src/framework/cl/cl_deleter.h b/src/framework/cl/cl_deleter.h new file mode 100644 index 0000000000000000000000000000000000000000..c18c51374deceb2e65afd897f2248b6ee627535d --- /dev/null +++ b/src/framework/cl/cl_deleter.h @@ -0,0 +1,52 @@ +/* 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. */ + +#pragma once + +#include + +struct CLKernelDeleter { + template + void operator()(T *clKernelObj) { + clReleaseKernel(clKernelObj); + } +}; + +struct CLMemDeleter { + template + void operator()(T *clMemObj) { + clReleaseMemObject(clMemObj); + } +}; + +struct CLCommQueueDeleter { + template + void operator()(T *clQueueObj) { + clReleaseCommandQueue(clQueueObj); + } +}; + +struct CLContextDeleter { + template + void operator()(T *clContextObj) { + clReleaseContext(clContextObj); + } +}; + +struct CLProgramDeleter { + template + void operator()(T *clProgramObj) { + clReleaseProgram(clProgramObj); + } +}; diff --git a/src/framework/cl/cl_engine.cpp b/src/framework/cl/cl_engine.cpp index 8cfa24c9c4a5afcb4027da89c311b79cd0a96e2e..6508e6c1626caab88d7dd652b04b34d38483fe1a 100644 --- a/src/framework/cl/cl_engine.cpp +++ b/src/framework/cl/cl_engine.cpp @@ -12,6 +12,7 @@ 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 "framework/cl/cl_tool.h" #include "framework/cl/cl_engine.h" #include @@ -25,11 +26,11 @@ bool CLEngine::Init() { cl_int status; setPlatform(); setClDeviceId(); - setClContext(); - setClCommandQueue(); - std::string filename = "./HelloWorld_Kernel.cl"; - loadKernelFromFile(filename.c_str()); - buildProgram(); +// setClContext(); +// setClCommandQueue(); +// std::string filename = "./HelloWorld_Kernel.cl"; +// loadKernelFromFile(filename.c_str()); +// buildProgram(); initialized_ = true; } @@ -38,19 +39,19 @@ CLEngine *CLEngine::Instance() { return &cl_engine_; } -std::unique_ptr<_cl_kernel, clKernel_deleter> CLEngine::GSetKernel( - const std::string &kernel_name) { - std::unique_ptr<_cl_kernel, clKernel_deleter> kernel( - clCreateKernel(program_.get(), kernel_name.c_str(), NULL)); - return std::move(kernel); -} - -bool CLEngine::SetClCommandQueue() { - cl_int status; - command_queue_.reset( - clCreateCommandQueue(context_.get(), devices_[0], 0, &status)); - return true; -} +//std::unique_ptr<_cl_kernel, clKernel_deleter> CLEngine::GSetKernel( +// const std::string &kernel_name) { +// std::unique_ptr<_cl_kernel, clKernel_deleter> kernel( +// clCreateKernel(program_.get(), kernel_name.c_str(), NULL)); +// return std::move(kernel); +//} +// +//bool CLEngine::SetClCommandQueue() { +// cl_int status; +// command_queue_.reset( +// clCreateCommandQueue(context_.get(), devices_[0], 0, &status)); +// return true; +//} bool CLEngine::SetPlatform() { platform_ = NULL; // the chosen platform @@ -70,10 +71,10 @@ bool CLEngine::SetPlatform() { } } -bool CLEngine::SetClContext() { - context_.reset(clCreateContext(NULL, 1, devices_, NULL, NULL, NULL)); - return true; -} +//bool CLEngine::SetClContext() { +// context_.reset(clCreateContext(NULL, 1, devices_, NULL, NULL, NULL)); +// return true; +//} bool CLEngine::SetClDeviceId() { cl_uint numDevices = 0; @@ -92,40 +93,35 @@ bool CLEngine::SetClDeviceId() { return false; } -bool CLEngine::LoadKernelFromFile(const char *kernel_file) { - size_t size; - char *str; - std::fstream f(kernel_file, (std::fstream::in | std::fstream::binary)); - - if (!f.is_open()) { - return false; - } - - size_t fileSize; - f.seekg(0, std::fstream::end); - size = fileSize = (size_t)f.tellg(); - f.seekg(0, std::fstream::beg); - str = new char[size + 1]; - if (!str) { - f.close(); - return 0; - } +//bool CLEngine::LoadKernelFromFile(const char *kernel_file) { +// size_t size; +// char *str; +// std::fstream f(kernel_file, (std::fstream::in | std::fstream::binary)); +// +// if (!f.is_open()) { +// return false; +// } +// +// size_t fileSize; +// f.seekg(0, std::fstream::end); +// size = fileSize = (size_t)f.tellg(); +// f.seekg(0, std::fstream::beg); +// str = new char[size + 1]; +// if (!str) { +// f.close(); +// return 0; +// } +// +// f.read(str, fileSize); +// f.close(); +// str[size] = '\0'; +// const char *source = str; +// size_t sourceSize[] = {strlen(source)}; +// program_.reset( +// clCreateProgramWithSource(context_.get(), 1, &source, sourceSize, NULL)); +// return true; +//} - f.read(str, fileSize); - f.close(); - str[size] = '\0'; - const char *source = str; - size_t sourceSize[] = {strlen(source)}; - program_.reset( - clCreateProgramWithSource(context_.get(), 1, &source, sourceSize, NULL)); - return true; -} - -bool CLEngine::BuildProgram() { - cl_int status; - status = clBuildProgram(program_.get(), 0, 0, "-cl-fast-relaxed-math", 0, 0); - return true; -} } // namespace framework } // namespace paddle_mobile diff --git a/src/framework/cl/cl_engine.h b/src/framework/cl/cl_engine.h index 77d5eee2afedd875ccb2d6b9cf3a3124ca8e2be0..5943635a404442c9190a01756ea2f4737859a7a2 100644 --- a/src/framework/cl/cl_engine.h +++ b/src/framework/cl/cl_engine.h @@ -21,60 +21,63 @@ limitations under the License. */ #include #include +#include "common/enforce.h" +#include "framework/cl/cl_deleter.h" + namespace paddle_mobile { namespace framework { -struct CLContext {}; - -struct CLKernelDeleter { - template - void operator()(T *clKernelObj) { - clReleaseKernel(clKernelObj); - } -}; - -struct CLMemDeleter { - template - void operator()(T *clMemObj) { - clReleaseMemObject(clMemObj); - } -}; - -struct CLCommQueueDeleter { - template - void operator()(T *clQueueObj) { - clReleaseCommandQueue(clQueueObj); - } -}; - -struct CLContextDeleter { - template - void operator()(T *clContextObj) { - clReleaseContext(clContextObj); - } -}; - -struct CLProgramDeleter { - template - void operator()(T *clProgramObj) { - clReleaseProgram(clProgramObj); - } -}; - class CLEngine { public: static CLEngine *Instance(); bool Init(); - std::unique_ptr<_cl_kernel, clKernel_deleter> GetKernel( - const std::string &kernel_name); + std::unique_ptr<_cl_context, CLContextDeleter> CreateContext() { + cl_context c = clCreateContext(NULL, 1, devices_, NULL, NULL, NULL); + std::unique_ptr<_cl_context, CLContextDeleter> context_ptr(c); + return std::move(context_ptr); + } - const cl_context GetContext() { return context_.get(); } + std::unique_ptr<_cl_command_queue, CLContextDeleter> CreateClCommandQueue() { + cl_int status; + cl_command_queue = clCreateCommandQueue(context_.get(), devices_[0], 0, &status); + std::unique_ptr<_cl_command_queue, CLCommQueueDeleter> command_queue_ptr(cl_command_queue); + return std::move(command_queue_ptr); + } - const cl_program GetProgram() { return program_.get(); } + std::unique_ptr<_cl_program, CLProgramDeleter> CreateProgramWith(cl_context context, std::string file_name) { + const char *kernel_file = file_name.c_str(); + size_t size; + char *str; + std::fstream f(kernel_file, (std::fstream::in | std::fstream::binary)); + + PADDLE_MOBILE_ENFORCE(f.is_open(), " file open failed") + + size_t fileSize; + f.seekg(0, std::fstream::end); + size = fileSize = (size_t)f.tellg(); + f.seekg(0, std::fstream::beg); + str = new char[size+1]; + + PADDLE_MOBILE_ENFORCE(str != NULL, " str null") + + f.read(str, fileSize); + f.close(); + str[size] = '\0'; + const char *source = str; + size_t sourceSize[] = {strlen(source)}; + cl_program p = clCreateProgramWithSource(context, 1, &source, sourceSize, NULL); + std::unique_ptr<_cl_program, CLProgramDeleter> program_ptr(p); + return std::move(program_ptr); + } - const cl_command_queue GetCommandQueue() { return command_queue_.get(); } + bool CLEngine::BuildProgram(cl_program program) { + cl_int status; + status = clBuildProgram(program, 0, 0, "-cl-fast-relaxed-math", 0, 0); + CL_CHECK_ERRORS(status); + return true; + } private: CLEngine() { initialized_ = false; } @@ -83,20 +86,25 @@ class CLEngine { bool SetClDeviceId(); - bool SetClContext(); +// bool SetClContext(); - bool SetClCommandQueue(); +// bool SetClCommandQueue(); - bool LoadKernelFromFile(const char *kernel_file); +// bool LoadKernelFromFile(const char *kernel_file); - bool BuildProgram(); +// bool BuildProgram(); bool initialized_; + cl_platform_id platform_; + cl_device_id *devices_; + std::unique_ptr<_cl_context, CLContextDeleter> context_; + std::unique_ptr<_cl_command_queue, CLCommQueueDeleter> command_queue_; - std::unique_ptr<_cl_program, clProgram_deleter> program_; + + std::unique_ptr<_cl_program, CLProgramDeleter> program_; }; } // namespace framework diff --git a/src/framework/cl/cl_helper.h b/src/framework/cl/cl_helper.h new file mode 100644 index 0000000000000000000000000000000000000000..a151cc050cc1547f9a0575e621ffbb59e560e5e3 --- /dev/null +++ b/src/framework/cl/cl_helper.h @@ -0,0 +1,54 @@ +/* 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. */ + +#pragma once + +#include + +#include "framework/cl/cl_scope.h" +#include "framework/cl/cl_deleter.h" + +namespace paddle_mobile { +namespace framework { + +class CLHelper { + public: + + CLHelper(CLScope *scope): scope_(scope) { + } + + void AddKernel(const std::string &kernel_name, const std::string &file_name) { + auto kernel = scope_->GetKernel(kernel_name, file_name); + kernels.emplace_back(kernel); + } + + cl_kernel KernelAt(const int index) { + return kernels[index].get(); + } + + cl_command_queue CLCommandQueue() { + return scope_->CommandQueue(); + } + + cl_context CLContext() { + return scope_->Context(); + } + + private: + CLScope *scope_; + std::vector> kernels; +}; + +} +} diff --git a/src/framework/cl/cl_image.h b/src/framework/cl/cl_image.h new file mode 100644 index 0000000000000000000000000000000000000000..8d611612686766a21d6e7562ea8fa1965ade9e81 --- /dev/null +++ b/src/framework/cl/cl_image.h @@ -0,0 +1,46 @@ +/* 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. */ + +#pragma once + +#include "framework/ddim.h" +#include "framework/tensor.h" + +namespace paddle_mobile { +namespace framework { + +class CLImage { + public: + CLImage(cl_context context, float *tensorInput, DDim ddim) : tensorDims_(ddim), context_(context) { + + } + + const DDim &TensorDim(); + + private: + cl_mem cl_image_; + DDim tensorDims_; + cl_context context_; +}; + +void TensorToCLImage(Tensor *tensor, CLImage *image) { + +} + +void CLImageToTensor(CLImage *image, Tensor *tensor) { + +} + +} +} \ No newline at end of file diff --git a/src/framework/cl/cl_scope.h b/src/framework/cl/cl_scope.h new file mode 100644 index 0000000000000000000000000000000000000000..4ea86e5ee066d8355c5213132c24c0736ea8fc21 --- /dev/null +++ b/src/framework/cl/cl_scope.h @@ -0,0 +1,73 @@ +/* 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. */ + +#pragma once + +#include +#include +#include + +#include "framework/cl/cl_tool.h" +#include "framework/cl/cl_engine.h" +#include "framework/cl/cl_deleter.h" + +namespace paddle_mobile { +namespace framework { + +class CLScope { + public: + CLScope() { + CLEngine *engin = CLEngine::Instance(); + context_ = engin->CreateContext(); + command_queue_ = engin->CreateClCommandQueue(); + } + + cl_command_queue CommandQueue() { + return command_queue_.get(); + } + + std::unique_ptr<_cl_kernel, CLKernelDeleter> GetKernel(const std::string &kernel_name, const std::string &file_name) { + auto program = Program(file_name); + std::unique_ptr<_cl_kernel, CLKernelDeleter> kernel(clCreateKernel(program, kernel_name.c_str(), NULL)); + return std::move(kernel); + } + + cl_context Context() { + return context_.get(); + } + + cl_program Program(const std::string &file_name) { + auto it = programs_.find(file_name); + if (it != programs_.end()) { + return it->second.get(); + } + + auto program = CLEngine::Instance()->CreateProgramWith(context_.get(), file_name); + programs_[file_name] = program; + + + status_ = clBuildProgram(program, 0, 0, 0, 0, 0); + CL_CHECK_ERRORS(status_); + return program; + } + + private: + cl_int status_; + std::unique_ptr<_cl_context, CLContextDeleter> context_; + std::unique_ptr<_cl_command_queue, CLCommQueueDeleter> command_queue_; + std::unordered_map> programs_; +}; + +} +} diff --git a/src/framework/cl/cl_tool.h b/src/framework/cl/cl_tool.h new file mode 100644 index 0000000000000000000000000000000000000000..93a8a4ef4e6d0134596c57b0639bc2a8f1f9bff6 --- /dev/null +++ b/src/framework/cl/cl_tool.h @@ -0,0 +1,25 @@ +/* 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. */ + +#pragma once + +#define CL_CHECK_ERRORS(ERR) \ + if(ERR != CL_SUCCESS) \ + { \ + printf( \ + "OpenCL error with code %s happened in file %s at line %d. Exiting.\n", \ + opencl_error_to_str(ERR), __FILE__, __LINE__ \ + ); \ + } +#endif \ No newline at end of file diff --git a/src/framework/operator.h b/src/framework/operator.h index 5252ee65a2a80910500f4085bb92b80829f9e45b..56355b1f3e0bb901c80cff75e70c4c64858d1469 100644 --- a/src/framework/operator.h +++ b/src/framework/operator.h @@ -18,6 +18,7 @@ limitations under the License. */ #include #include + #include "common/enforce.h" #include "common/type_define.h" #include "common/types.h" @@ -31,6 +32,8 @@ limitations under the License. */ #include "framework/scope.h" #include "framework/tensor.h" #include "framework/variable.h" +#include "framework/cl/cl_scope.h" +#include "framework/cl/cl_helper.h" namespace paddle_mobile { namespace framework { @@ -112,7 +115,8 @@ class OperatorWithKernel : public OperatorBase { const VariableNameMap &outputs, const AttributeMap &attrs, std::shared_ptr scope) : OperatorBase(type, inputs, outputs, attrs, scope), - param_(inputs, outputs, attrs, *scope) {} + param_(inputs, outputs, attrs, *scope), + kernel_(scope->GetCLScpoe()) {} virtual void RunImpl() const { this->kernel_.Compute(this->param_); } @@ -138,6 +142,11 @@ class OperatorWithKernel : public OperatorBase { template class OpKernelBase { public: + OpKernelBase() = default; + + OpKernelBase(CLScope *clscope): cl_helper_(clscope) { + } + /* * @b 所有kernel 需实现 Compute 方法 * @p para 这个参数为 kernel 运算时所需要用到参数组成的一个结构体, @@ -158,6 +167,10 @@ class OpKernelBase { #ifdef PADDLE_MOBILE_MALI_GPU void *acl_op_; #endif + + CLHelper cl_helper_; + + }; #define DEFINE_OP_CONSTRUCTOR(cls, parent_cls) \ diff --git a/src/framework/scope.h b/src/framework/scope.h index 054f141ff68895e0879fd31e15d90c76ea038135..06ae469cc5c7cecf19d81c8fcb18b79ce7efc23d 100644 --- a/src/framework/scope.h +++ b/src/framework/scope.h @@ -15,6 +15,11 @@ limitations under the License. */ #pragma once #include + +//#ifdef PADDLE_MOBILE_OCL +#include "framework/cl/cl_scope.h" +//#endif + #include #include "variable.h" @@ -33,6 +38,11 @@ class Scope { delete kid; } kids_.clear(); + +//#ifdef PADDLE_MOBILE_OCL + delete cl_scope_; +//#endif + } Scope &NewScope() const; @@ -72,6 +82,10 @@ class Scope { Variable *FindVarLocally(const std::string &name) const; + CLScope *GetCLScpoe() { + return cl_scope_; + } + private: // Call Scope::NewScope for a sub-scope. explicit Scope(Scope const *parent) : parent_(parent) {} @@ -79,6 +93,9 @@ class Scope { mutable std::unordered_map vars_; mutable std::list kids_; Scope const *parent_{nullptr}; +//#ifdef PADDLE_MOBILE_OCL + CLScope *cl_scope_ = new CLScope(); +//#endif }; } // namespace framework } // namespace paddle_mobile diff --git a/src/io/executor.cpp b/src/io/executor.cpp index 33a6ff359515b0cb6f8e9c2dd2c10af6001490e5..926917d178569525f8bf878748dc0b8085ac1152 100644 --- a/src/io/executor.cpp +++ b/src/io/executor.cpp @@ -290,6 +290,7 @@ void Executor::InitCombineMemory() { delete origin_data; LOG(kLOG_INFO) << " end init combine memory "; } + template bool Executor::varInputMemory( const std::shared_ptr &var_desc, Variable *var, diff --git a/src/operators/kernel/cl/cl_kernel/conv_kernel.cl b/src/operators/kernel/cl/cl_kernel/conv_kernel.cl new file mode 100644 index 0000000000000000000000000000000000000000..71bd1d9ceec4091276d9143d7ad1913371ccbad1 --- /dev/null +++ b/src/operators/kernel/cl/cl_kernel/conv_kernel.cl @@ -0,0 +1,7 @@ + + + +__kernel void conv_3x3(__global float* in, __global float* out) { + int num = get_global_id(0); + out[num] = in[num] * 0.1 + 102; + } diff --git a/src/operators/kernel/cl/conv_kernel.cpp b/src/operators/kernel/cl/conv_kernel.cpp new file mode 100644 index 0000000000000000000000000000000000000000..6e5a4326261f1ad17d0c803ee3a4cfb4156eccf0 --- /dev/null +++ b/src/operators/kernel/cl/conv_kernel.cpp @@ -0,0 +1,39 @@ +/* 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 CONV_OP + +#include "operators/kernel/conv_kernel.h" +#include "operators/kernel/central-arm-func/conv_arm_func.h" + +namespace paddle_mobile { +namespace operators { + +template <> +bool ConvKernel::Init(ConvParam *param) { + + return true; +} + +template <> +void ConvKernel::Compute(const ConvParam ¶m) const { + +} + +template class ConvKernel; + +} // namespace operators +} // namespace paddle_mobile + +#endif