m53297601
/
mindspore
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

!4506 [MS][LITE][Develop]add caffe prelu in opencl

Browse Source 
Merge pull request !4506 from liuzhongkai/caffe_prelu
pull/4506/MERGE
mindspore-ci-bot 5 years ago committed by Gitee
parent abe89923d8 89003fd5f4
commit bf8bdd6128
9 changed files with 390 additions and 5 deletions
Show Stats Download Patch File Download Diff File

2
mindspore/lite/src/runtime/kernel/opencl/CMakeLists.txt
Unescape View File

@ -11,4 +11,6 @@ set(OPENCL_KERNEL_SRC
${CMAKE_CURRENT_SOURCE_DIR}/kernel/conv2d_transpose.cc
${CMAKE_CURRENT_SOURCE_DIR}/kernel/transpose.cc
${CMAKE_CURRENT_SOURCE_DIR}/kernel/reshape.cc
${CMAKE_CURRENT_SOURCE_DIR}/kernel/activation.cc
${CMAKE_CURRENT_SOURCE_DIR}/kernel/caffe_prelu.cc
)

26
mindspore/lite/src/runtime/kernel/opencl/cl/fp32/caffe_prelu.cl
Unescape View File

@ -0,0 +1,26 @@
#pragma OPENCL EXTENSION cl_arm_printf : enable
#define SLICES 4
#define UP_DIV(x, y) (((x) + (y) - (1)) / (y))
#define FLT4 float4
#define READ_FLT4 read_imagef
#define WRITE_FLT4 write_imagef
__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
__kernel void CaffePRelu(__read_only image2d_t input, __write_only image2d_t output, const int4 input_shape,
__global float *alpha) {
int C = input_shape.w; // channel size
int Y = get_global_id(0); // height id
int X = get_global_id(1); // weight id
for (int num = 0; num < UP_DIV(C, SLICES); ++num) {
FLT4 in_c4 = READ_FLT4(input, smp_zero, (int2)(X * UP_DIV(C, SLICES) + num, Y)); // NHWC4: H WC
FLT4 tmp;
int index = num * 4;
tmp.x = in_c4.x * alpha[index];
tmp.y = in_c4.y * alpha[index + 1];
tmp.z = in_c4.z * alpha[index + 2];
tmp.w = in_c4.w * alpha[index + 3];
WRITE_FLT4(output, (int2)(X * UP_DIV(C, SLICES) + num, Y), tmp); // NHWC4: H WC
}
}

3
mindspore/lite/src/runtime/kernel/opencl/kernel/activation.cc
Unescape View File

@ -44,7 +44,6 @@ int ActivationOpenClKernel::Init() {
MS_LOG(ERROR) << "Activate fun only support dim=4, but your dim=" << in_tensors_[0]->shape().size();
return RET_ERROR;
}
ori_format_ = out_tensors_[0]->GetFormat();
std::string program_name = "";
std::string kernel_name = "";
std::string source = activation_source_fp32;
@ -68,6 +67,8 @@ int ActivationOpenClKernel::Init() {
auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance();
ocl_runtime->LoadSource(program_name, source);
ocl_runtime->BuildKernel(kernel_, program_name, kernel_name, build_options);
ori_format_ = out_tensors_[0]->GetFormat();
out_tensors_[0]->SetFormat(schema::Format_NHWC4);
MS_LOG(DEBUG) << op_parameter_->name_ << " init Done!";
return RET_OK;
}

150
mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.cc
Unescape View File

@ -0,0 +1,150 @@
/**
* 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.
*/
#include <string>
#include <set>
#include <vector>
#include "src/kernel_registry.h"
#include "include/errorcode.h"
#include "src/runtime/kernel/opencl/kernel/caffe_prelu.h"
#include "src/runtime/opencl/opencl_runtime.h"
#include "src/runtime/kernel/opencl/cl/fp32/caffe_prelu.cl.inc"
#include "src/runtime/kernel/arm/nnacl/caffeprelu.h"
using mindspore::kernel::KERNEL_ARCH::kGPU;
using mindspore::lite::KernelRegistrar;
using mindspore::lite::RET_ERROR;
using mindspore::lite::RET_OK;
using mindspore::schema::PrimitiveType_CaffePReLU;
namespace mindspore::kernel {
void CaffePReluOpenCLKernel::CaffeWeight() {
int C = in_tensors_[1]->shape()[0];
int div_ci = UP_DIV(C, C4NUM);
std::cout << div_ci << std::endl;
auto allocator = lite::opencl::OpenCLRuntime::GetInstance()->GetAllocator();
CaffeWeight_ = reinterpret_cast<FLOAT_t *>(allocator->Malloc(div_ci * C4NUM * sizeof(FLOAT_t)));
CaffeWeight_ = reinterpret_cast<FLOAT_t *>(allocator->MapBuffer(CaffeWeight_, CL_MAP_WRITE, nullptr, true));
memset(CaffeWeight_, 0x00, div_ci * C4NUM * sizeof(FLOAT_t));
auto origin_weight = reinterpret_cast<FLOAT_t *>(in_tensors_[1]->Data());
for (int i = 0; i < in_tensors_[1]->ElementsNum(); ++i) {
CaffeWeight_[i] = origin_weight[i];
}
allocator->UnmapBuffer(CaffeWeight_);
}
int CaffePReluOpenCLKernel::Init() {
if (in_tensors_[0]->shape().size() != 4) {
MS_LOG(ERROR) << "Caffe PRelu only support dim=4, but your dim=" << in_tensors_[0]->shape().size();
return RET_ERROR;
}
CaffeWeight();
std::set<std::string> build_options;
std::string source = caffe_prelu_source_fp32;
std::string program_name = "CaffePRelu";
std::string kernel_name = "CaffePRelu";
auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance();
ocl_runtime->LoadSource(program_name, source);
ocl_runtime->BuildKernel(kernel_, program_name, kernel_name, build_options);
ori_format_ = out_tensors_[0]->GetFormat();
out_tensors_[0]->SetFormat(schema::Format_NHWC4);
MS_LOG(DEBUG) << program_name << " Init Done!";
return RET_OK;
}
int CaffePReluOpenCLKernel::Run() {
int N = in_tensors_[0]->shape()[0];
int H = in_tensors_[0]->shape()[1];
int W = in_tensors_[0]->shape()[2];
int C = in_tensors_[0]->shape()[3];
cl_int4 input_shape = {N, H, W, C};
int C_Weight = in_tensors_[1]->shape()[0];
if (UP_DIV(C_Weight, C4NUM) != UP_DIV(C, C4NUM)) {
MS_LOG(ERROR) << "CaffePRelu weight channel size:" << C_Weight
<< " must be equal with in_teneors channel size:" << C;
return RET_ERROR;
}
MS_LOG(DEBUG) << op_parameter_->name_ << " Running!";
auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance();
int arg_idx = 0;
ocl_runtime->SetKernelArg(kernel_, arg_idx++, in_tensors_[0]->Data());
ocl_runtime->SetKernelArg(kernel_, arg_idx++, out_tensors_[0]->Data());
ocl_runtime->SetKernelArg(kernel_, arg_idx++, input_shape);
ocl_runtime->SetKernelArg(kernel_, arg_idx++, CaffeWeight_);
std::vector<size_t> local = {1, 1};
std::vector<size_t> global = {static_cast<size_t>(H), static_cast<size_t>(W)};
auto ret = ocl_runtime->RunKernel(kernel_, global, local, nullptr);
if (ret != RET_OK) {
MS_LOG(ERROR) << "Run kernel " << op_parameter_->name_ << " error.";
return RET_ERROR;
}
return RET_OK;
}
int CaffePReluOpenCLKernel::GetImageSize(size_t idx, std::vector<size_t> *img_size) {
int H = in_tensors_[0]->shape()[1];
int W = in_tensors_[0]->shape()[2];
int C = in_tensors_[0]->shape()[3];
#ifdef ENABLE_FP16
size_t img_dtype = CL_HALF_FLOAT;
#else
size_t img_dtype = CL_FLOAT;
#endif
img_size->clear();
img_size->push_back(W * UP_DIV(C, C4NUM));
img_size->push_back(H);
img_size->push_back(img_dtype);
return RET_OK;
}
kernel::LiteKernel *OpenCLCaffePReluKernelCreator(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) {
if (inputs.size() == 0) {
MS_LOG(ERROR) << "Input data size must be greater than 0, but your size is " << inputs.size();
return nullptr;
}
if (inputs[0]->shape()[0] > 1) {
MS_LOG(ERROR) << "Init CaffePRelu kernel failed: Unsupported multi-batch.";
return nullptr;
}
auto *kernel =
new (std::nothrow) CaffePReluOpenCLKernel(reinterpret_cast<OpParameter *>(opParameter), inputs, outputs);
if (kernel == nullptr) {
MS_LOG(ERROR) << "Kernel " << opParameter->name_ << "is nullptr.";
return nullptr;
}
auto ret = kernel->Init();
if (ret != RET_OK) {
MS_LOG(ERROR) << "Init CaffePRelu kernel failed!";
delete kernel;
return nullptr;
}
return kernel;
}
REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_CaffePReLU, OpenCLCaffePReluKernelCreator)
} // namespace mindspore::kernel

49
mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.h
Unescape View File

@ -0,0 +1,49 @@
/**
* 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_RUNTIME_KERNEL_OPENCL_KERNEL_CAFFEPRELU_H_
#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_OPENCL_KERNEL_CAFFEPRELU_H_
#include <vector>
#include <string>
#include "src/ir/tensor.h"
#include "src/runtime/kernel/opencl/opencl_kernel.h"
#include "schema/model_generated.h"
#include "src/runtime/opencl/opencl_runtime.h"
namespace mindspore::kernel {
class CaffePReluOpenCLKernel : public OpenCLKernel {
public:
explicit CaffePReluOpenCLKernel(OpParameter *parameter, const std::vector<lite::tensor::Tensor *> &inputs,
const std::vector<lite::tensor::Tensor *> &outputs)
: OpenCLKernel(parameter, inputs, outputs) {}
~CaffePReluOpenCLKernel() override{};
int Init() override;
int Run() override;
int GetImageSize(size_t idx, std::vector<size_t> *img_size) override;
void CaffeWeight();
private:
cl::Kernel kernel_;
FLOAT_t *CaffeWeight_;
};
} // namespace mindspore::kernel
#endif // MINDSPORE_LITE_SRC_RUNTIME_KERNEL_OPENCL_KERNEL_CAFFEPRELU_H_

3
mindspore/lite/src/runtime/opencl/opencl_allocator.cc
Unescape View File

@ -122,7 +122,8 @@ void *OpenCLAllocator::Malloc(size_t size, const std::vector<size_t> &img_size)
mem_buf->host_ptr_ = host_ptr;
mem_buf->image_ptr_ = image_ptr;
mem_buf->img_size = img_size;
MS_LOG(DEBUG) << "Malloc a new buffer. size: " << mem_buf->size_ << ", host addr: " << mem_buf->host_ptr_
std::string type_name = img_size.empty() ? "buffer" : "Image2D";
MS_LOG(DEBUG) << "Malloc a new " << type_name << ". size: " << mem_buf->size_ << ", host addr: " << mem_buf->host_ptr_
<< ", device addr: " << mem_buf->device_ptr_ << ", image_addr: " << image_ptr;
allocated_list_[host_ptr] = mem_buf.release();
UnLock();

5
mindspore/lite/test/CMakeLists.txt
Unescape View File

@ -149,6 +149,7 @@ if (SUPPORT_GPU)
${LITE_DIR}/src/runtime/kernel/opencl/kernel/transpose.cc
${LITE_DIR}/src/runtime/kernel/opencl/kernel/reshape.cc
${LITE_DIR}/src/runtime/kernel/opencl/kernel/to_format.cc
${LITE_DIR}/src/runtime/kernel/opencl/kernel/caffe_prelu.cc
)
endif()
### minddata lite
@ -322,15 +323,15 @@ if (SUPPORT_GPU)
${TEST_DIR}/ut/src/runtime/kernel/opencl/transpose_tests.cc
${TEST_DIR}/ut/src/runtime/kernel/opencl/convolution_tests.cc
${TEST_DIR}/ut/src/runtime/kernel/opencl/activation_tests.cc
#${TEST_DIR}/ut/src/runtime/kernel/opencl/leakyrelu_tests.cc
${TEST_DIR}/ut/src/runtime/kernel/opencl/to_format_tests.cc
${TEST_DIR}/ut/src/runtime/kernel/opencl/caffe_prelu_tests.cc
)
endif()
if (ENABLE_FP16)
set(TEST_SRC
${TEST_SRC}
)
${TEST_DIR}/ut/src/runtime/kernel/arm/fp16/convolution_fp16_tests.cc)
endif ()

2
mindspore/lite/test/ut/src/runtime/kernel/opencl/activation_tests.cc
Unescape View File

@ -128,7 +128,7 @@ int RunSubGraphOpenCLKernel(const std::vector<lite::tensor::Tensor *> &inputs,
return RET_OK;
}
TEST_F(TestActivationOpenCL, LeakyReluFp32_dim4) {
TEST_F(TestActivationOpenCL, ActivationFp32_dim4) {
MS_LOG(INFO) << "Begin test:";
auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance();
ocl_runtime->Init();

155
mindspore/lite/test/ut/src/runtime/kernel/opencl/caffe_prelu_tests.cc
Unescape View File

@ -0,0 +1,155 @@
/**
* 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.
*/
#include <iostream>
#include "utils/log_adapter.h"
#include "common/common_test.h"
#include "mindspore/lite/src/common/file_utils.h"
#include "src/runtime/kernel/arm/nnacl/pack.h"
#include "mindspore/lite/src/runtime/opencl/opencl_runtime.h"
#include "mindspore/lite/src/runtime/kernel/opencl/subgraph_opencl_kernel.h"
#include "mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.h"
#include "mindspore/lite/src/runtime/kernel/arm/nnacl/caffeprelu.h"
using mindspore::kernel::CaffePReluOpenCLKernel;
using mindspore::kernel::LiteKernel;
using mindspore::kernel::SubGraphOpenCLKernel;
using mindspore::lite::RET_ERROR;
using mindspore::lite::RET_OK;
namespace mindspore {
class TestCaffePReluOpenCL : public mindspore::CommonTest {};
void LoadDataCaffePRelu(void *dst, size_t dst_size, const std::string &file_path) {
if (file_path.empty()) {
memset(dst, 0x00, dst_size);
} else {
auto src_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(file_path.c_str(), &dst_size));
memcpy(dst, src_data, dst_size);
}
}
void CompareOutCaffePRelu(lite::tensor::Tensor *output_tensor, const std::string &standard_answer_file) {
auto *output_data = reinterpret_cast<float *>(output_tensor->Data());
size_t output_size = output_tensor->ElementsC4Num();
auto expect_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(standard_answer_file.c_str(), &output_size));
constexpr float atol = 0.0002;
for (int i = 0; i < output_tensor->ElementsC4Num(); ++i) {
if (std::fabs(output_data[i] - expect_data[i]) > atol) {
printf("error at idx[%d] expect=%.3f output=%.3f\n", i, expect_data[i], output_data[i]);
printf("error at idx[%d] expect=%.3f output=%.3f\n", i, expect_data[i], output_data[i]);
printf("error at idx[%d] expect=%.3f output=%.3f\n\n\n", i, expect_data[i], output_data[i]);
return;
}
}
printf("compare success!\n");
printf("compare success!\n");
printf("compare success!\n\n\n");
}
void printf_tensor_caffeprelu(mindspore::lite::tensor::Tensor *in_data, int size) {
auto input_data = reinterpret_cast<float *>(in_data->Data());
for (int i = 0; i < size; ++i) {
printf("%f ", input_data[i]);
}
printf("\n");
MS_LOG(INFO) << "Print tensor done";
}
void printf_float(float *data, int num = 0) {
float *temp = data;
for (int i = 0; i < num; ++i) {
std::cout << *temp << " ";
temp++;
}
std::cout << std::endl;
}
TEST_F(TestCaffePReluOpenCL, CaffePReluFp32_dim4) {
std::string in_file = "/data/local/tmp/in_data.bin";
std::string weight_file = "/data/local/tmp/weight_data.bin";
std::string standard_answer_file = "/data/local/tmp/caffeprelu.bin";
MS_LOG(INFO) << "CaffePRelu Begin test:";
auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance();
ocl_runtime->Init();
auto allocator = ocl_runtime->GetAllocator();
MS_LOG(INFO) << "CaffePRelu init tensors.";
std::vector<int> input_shape = {1, 4, 3, 9};
std::vector<int> output_shape = {1, 4, 3, 9};
auto data_type = kNumberTypeFloat32;
auto tensor_type = schema::NodeType_ValueNode;
auto *input_tensor = new lite::tensor::Tensor(data_type, input_shape, schema::Format_NHWC, tensor_type);
auto *output_tensor = new lite::tensor::Tensor(data_type, output_shape, schema::Format_NHWC4, tensor_type);
auto *weight_tensor =
new lite::tensor::Tensor(data_type, std::vector<int>{input_shape[3]}, schema::Format_NHWC, tensor_type);
std::vector<lite::tensor::Tensor *> inputs{input_tensor, weight_tensor};
std::vector<lite::tensor::Tensor *> outputs{output_tensor};
std::cout << input_tensor->ElementsNum() << std::endl;
std::cout << input_tensor->ElementsC4Num() << std::endl;
// freamework to do!!! allocate memory by hand
inputs[0]->MallocData(allocator);
inputs[1]->MallocData(allocator);
std::cout << input_tensor->Size() << std::endl;
LoadDataCaffePRelu(input_tensor->Data(), input_tensor->Size(), in_file);
MS_LOG(INFO) << "CaffePRelu==================input data================";
printf_tensor_caffeprelu(inputs[0], input_tensor->ElementsNum());
LoadDataCaffePRelu(weight_tensor->Data(), weight_tensor->Size(), weight_file);
MS_LOG(INFO) << "CaffePRelu==================weight data================";
printf_tensor_caffeprelu(inputs[1], weight_tensor->ElementsNum());
auto param = new CaffePReluParameter();
param->channel_num_ = input_shape[3];
auto *caffeprelu_kernel =
new (std::nothrow) kernel::CaffePReluOpenCLKernel(reinterpret_cast<OpParameter *>(param), inputs, outputs);
if (caffeprelu_kernel == nullptr) {
MS_LOG(ERROR) << "Create caffe prelu kernel error.";
return;
}
auto ret = caffeprelu_kernel->Init();
if (ret != RET_OK) {
MS_LOG(ERROR) << "caffeprelu_kernel init error.";
return;
}
MS_LOG(INFO) << "initialize sub_graph";
std::vector<kernel::LiteKernel *> kernels{caffeprelu_kernel};
auto *sub_graph = new (std::nothrow) kernel::SubGraphOpenCLKernel({input_tensor}, outputs, kernels, kernels, kernels);
if (sub_graph == nullptr) {
MS_LOG(ERROR) << "Create sub_graph kernel error.";
return;
}
ret = sub_graph->Init();
if (ret != RET_OK) {
MS_LOG(ERROR) << "sub_graph init error.";
return;
}
MS_LOG(INFO) << "Sub graph begin running!";
ret = sub_graph->Run();
if (ret != RET_OK) {
MS_LOG(ERROR) << "sub_graph run error.";
return;
}
MS_LOG(INFO) << "CaffePRelu==================output data================";
printf_tensor_caffeprelu(outputs[0], output_tensor->ElementsC4Num());
std::cout << "output date size:" << output_tensor->Size() << std::endl;
CompareOutCaffePRelu(output_tensor, standard_answer_file);
}
} // namespace mindspore
Write Preview
Loading…
Cancel
Save