diff --git a/mindspore/lite/src/runtime/kernel/opencl/cl/caffe_prelu.cl b/mindspore/lite/src/runtime/kernel/opencl/cl/caffe_prelu.cl deleted file mode 100644 index bfaec1dfe1fbba14f6e746f5f8b92744c66a41c2..0000000000000000000000000000000000000000 --- a/mindspore/lite/src/runtime/kernel/opencl/cl/caffe_prelu.cl +++ /dev/null @@ -1,23 +0,0 @@ -#pragma OPENCL EXTENSION cl_arm_printf : enable - -#define SLICES 4 -#define UP_DIV(x, y) (((x) + (y) - (1)) / (y)) -__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_IMAGE(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_IMAGE(output, (int2)(X * UP_DIV(C, SLICES) + num, Y), tmp); // NHWC4: H WC - } -} diff --git a/mindspore/lite/src/runtime/kernel/opencl/cl/prelu.cl b/mindspore/lite/src/runtime/kernel/opencl/cl/prelu.cl new file mode 100644 index 0000000000000000000000000000000000000000..40fbb4cfe3fef31540f7400aec5d908d2f24f8e5 --- /dev/null +++ b/mindspore/lite/src/runtime/kernel/opencl/cl/prelu.cl @@ -0,0 +1,30 @@ +#pragma OPENCL EXTENSION cl_arm_printf : enable + +#define SLICES 4 +#define UP_DIV(x, y) (((x) + (y) - (1)) / (y)) +__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; + +__kernel void PRelu(__read_only image2d_t input, __write_only image2d_t output, const int4 input_shape, + __global float *alpha, const int dim) { + 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_IMAGE(input, smp_zero, (int2)(X * UP_DIV(C, SLICES) + num, Y)); // NHWC4: H WC + FLT4 tmp; + if (dim == 1) { + tmp.x = in_c4.x >= 0 ? in_c4.x : in_c4.x * (*alpha); + tmp.y = in_c4.y >= 0 ? in_c4.y : in_c4.y * (*alpha); + tmp.z = in_c4.z >= 0 ? in_c4.z : in_c4.z * (*alpha); + tmp.w = in_c4.w >= 0 ? in_c4.w : in_c4.w * (*alpha); + } else { + int index = num * 4; + tmp.x = in_c4.x >= 0 ? in_c4.x : in_c4.x * alpha[index]; + tmp.y = in_c4.y >= 0 ? in_c4.y : in_c4.y * alpha[index + 1]; + tmp.z = in_c4.z >= 0 ? in_c4.z : in_c4.z * alpha[index + 2]; + tmp.w = in_c4.w >= 0 ? in_c4.w : in_c4.w * alpha[index + 3]; + } + WRITE_IMAGE(output, (int2)(X * UP_DIV(C, SLICES) + num, Y), tmp); // NHWC4: H WC + } +} diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.cc deleted file mode 100644 index 07a2d517548104c6c8ad08c007cab88249374ddc..0000000000000000000000000000000000000000 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.cc +++ /dev/null @@ -1,152 +0,0 @@ -/** - * 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 -#include -#include - -#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/caffe_prelu.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_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(allocator->Malloc(div_ci * C4NUM * sizeof(FLOAT_t))); - CaffeWeight_ = reinterpret_cast(allocator->MapBuffer(CaffeWeight_, CL_MAP_WRITE, nullptr, true)); - memset(CaffeWeight_, 0x00, div_ci * C4NUM * sizeof(FLOAT_t)); - auto origin_weight = reinterpret_cast(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 build_options; - std::string source = caffe_prelu_source; - 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); - in_ori_format_ = in_tensors_[0]->GetFormat(); - in_tensors_[0]->SetFormat(schema::Format_NHWC4); - out_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 local = {1, 1}; - std::vector global = {static_cast(H), static_cast(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 *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 &inputs, - const std::vector &outputs, - OpParameter *opParameter, const lite::Context *ctx, - const kernel::KernelKey &desc, - const mindspore::lite::PrimitiveC *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), 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 diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.h b/mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.h deleted file mode 100644 index 9a25aef4fa5ae32c9bb8913cd364267b3cf30849..0000000000000000000000000000000000000000 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/caffe_prelu.h +++ /dev/null @@ -1,49 +0,0 @@ -/** - * 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 -#include - -#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 &inputs, - const std::vector &outputs) - : OpenCLKernel(parameter, inputs, outputs) {} - ~CaffePReluOpenCLKernel() override{}; - - int Init() override; - int Run() override; - int GetImageSize(size_t idx, std::vector *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_ diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc index 4f4e39a63af2e2f91f0ba1bda7793487795d36e2..9f398a958e6f4cac8b480a831a1659dec3bc4f11 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc @@ -23,8 +23,7 @@ #include "include/errorcode.h" #include "src/runtime/kernel/opencl/kernel/prelu.h" #include "src/runtime/opencl/opencl_runtime.h" -#include "src/runtime/kernel/opencl/cl/activation.cl.inc" -#include "nnacl/prelu_parameter.h" +#include "src/runtime/kernel/opencl/cl/prelu.cl.inc" using mindspore::kernel::KERNEL_ARCH::kGPU; using mindspore::lite::KernelRegistrar; @@ -34,15 +33,41 @@ using mindspore::schema::PrimitiveType_Prelu; namespace mindspore::kernel { +void PReluOpenCLKernel::InitBuffer() { + 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(); + PReluWeight_ = reinterpret_cast(allocator->Malloc(div_ci * C4NUM * sizeof(FLOAT_t))); + PReluWeight_ = reinterpret_cast(allocator->MapBuffer(PReluWeight_, CL_MAP_WRITE, nullptr, true)); + memset(PReluWeight_, 0x00, div_ci * C4NUM * sizeof(FLOAT_t)); + auto origin_weight = reinterpret_cast(in_tensors_[1]->Data()); + for (int i = 0; i < in_tensors_[1]->ElementsNum(); ++i) { + PReluWeight_[i] = origin_weight[i]; + } + allocator->UnmapBuffer(PReluWeight_); +} + int PReluOpenCLKernel::Init() { if (in_tensors_[0]->shape().size() != 4) { MS_LOG(ERROR) << "PRelu only support dim=4, but your dim=" << in_tensors_[0]->shape().size(); return RET_ERROR; } + int C_Weight = in_tensors_[1]->shape()[0]; + int C = in_tensors_[0]->shape()[3]; + if (C_Weight != 1 && UP_DIV(C_Weight, C4NUM) != UP_DIV(C, C4NUM)) { + MS_LOG(ERROR) + << "PRelu weight channel size must be 1 or must be equal with in_teneors channel size, but your weight size is " + << C_Weight << " and your input channel size is " << C; + return RET_ERROR; + } + if (C_Weight != 1) { + InitBuffer(); + } std::set build_options; - std::string source = activation_source; + std::string source = prelu_source; std::string program_name = "PRelu"; - std::string kernel_name = "ReluScalar"; + std::string kernel_name = "PRelu"; auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance(); ocl_runtime->LoadSource(program_name, source); ocl_runtime->BuildKernel(kernel_, program_name, kernel_name, build_options); @@ -61,17 +86,18 @@ int PReluOpenCLKernel::Run() { int W = in_tensors_[0]->shape()[2]; int C = in_tensors_[0]->shape()[3]; cl_int4 input_shape = {N, H, W, C}; - if (in_tensors_[1]->ElementsNum() < 1) { - MS_LOG(ERROR) << "PRelu weight size must be greater than 1! But your weight size is " - << in_tensors_[1]->ElementsNum(); - return RET_ERROR; - } + 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++, reinterpret_cast(in_tensors_[1]->Data())[0]); + if (in_tensors_[1]->shape()[0] == 1) { + ocl_runtime->SetKernelArg(kernel_, arg_idx++, reinterpret_cast(in_tensors_[1]->Data())); + } else { + ocl_runtime->SetKernelArg(kernel_, arg_idx++, PReluWeight_); + } + ocl_runtime->SetKernelArg(kernel_, arg_idx++, reinterpret_cast(in_tensors_[1]->shape()[0])); std::vector local = {1, 1}; std::vector global = {static_cast(H), static_cast(W)}; diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.h b/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.h index 79b232ce32239891b079bbe2a4dcbda94069b252..cc2429250a9d5344bf5d3ca043f3e7021f6e6a5b 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.h +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.h @@ -36,9 +36,11 @@ class PReluOpenCLKernel : public OpenCLKernel { int Init() override; int Run() override; int GetImageSize(size_t idx, std::vector *img_size) override; + void InitBuffer(); private: cl::Kernel kernel_; + FLOAT_t *PReluWeight_; }; } // namespace mindspore::kernel diff --git a/mindspore/lite/test/CMakeLists.txt b/mindspore/lite/test/CMakeLists.txt index 6767810eb84a96d4041778a83f0545e4dfe9103f..51aa509a00d002f0b47a659458648450edc2d7f0 100644 --- a/mindspore/lite/test/CMakeLists.txt +++ b/mindspore/lite/test/CMakeLists.txt @@ -155,7 +155,6 @@ 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 ${LITE_DIR}/src/runtime/kernel/opencl/kernel/prelu.cc ${LITE_DIR}/src/runtime/kernel/opencl/kernel/to_format.cc ${LITE_DIR}/src/runtime/kernel/opencl/kernel/biasadd.cc @@ -336,7 +335,6 @@ if (SUPPORT_GPU) ${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/to_format_tests.cc - ${TEST_DIR}/ut/src/runtime/kernel/opencl/caffe_prelu_tests.cc ${TEST_DIR}/ut/src/runtime/kernel/opencl/prelu_tests.cc ${TEST_DIR}/ut/src/runtime/kernel/opencl/reshape_tests.cc ${TEST_DIR}/ut/src/runtime/kernel/opencl/biasadd_tests.cc diff --git a/mindspore/lite/test/ut/src/runtime/kernel/opencl/caffe_prelu_tests.cc b/mindspore/lite/test/ut/src/runtime/kernel/opencl/caffe_prelu_tests.cc deleted file mode 100644 index ac6591be4ecb3947e2d1c68a45088ef5e9d64ab0..0000000000000000000000000000000000000000 --- a/mindspore/lite/test/ut/src/runtime/kernel/opencl/caffe_prelu_tests.cc +++ /dev/null @@ -1,203 +0,0 @@ -/** - * 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 - -#include "utils/log_adapter.h" -#include "common/common_test.h" -#include "mindspore/lite/src/common/file_utils.h" -#include "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/nnacl/prelu_parameter.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(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(output_tensor->Data()); - size_t output_size = output_tensor->ElementsC4Num(); - auto expect_data = reinterpret_cast(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(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 input_shape = {1, 4, 3, 9}; - std::vector output_shape = {1, 4, 3, 9}; - auto data_type = kNumberTypeFloat32; - auto tensor_type = schema::NodeType_ValueNode; - auto *input_tensor = - new (std::nothrow) lite::tensor::Tensor(data_type, input_shape, schema::Format_NHWC, tensor_type); - if (input_tensor == nullptr) { - MS_LOG(ERROR) << "new input tensor error"; - return; - } - auto *output_tensor = new lite::tensor::Tensor(data_type, output_shape, schema::Format_NHWC4, tensor_type); - if (output_tensor == nullptr) { - MS_LOG(ERROR) << "new output_tensor error"; - delete input_tensor; - return; - } - auto *weight_tensor = new (std::nothrow) - lite::tensor::Tensor(data_type, std::vector{input_shape[3]}, schema::Format_NHWC, tensor_type); - if (weight_tensor == nullptr) { - MS_LOG(ERROR) << "new weight_tensor error"; - delete input_tensor; - delete output_tensor; - return; - } - - std::vector inputs{input_tensor, weight_tensor}; - std::vector outputs{output_tensor}; - 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 (std::nothrow) PReluParameter(); - if (param == nullptr) { - MS_LOG(ERROR) << "new param error!"; - delete input_tensor; - delete output_tensor; - delete weight_tensor; - return; - } - param->channel_num_ = input_shape[3]; - auto *caffeprelu_kernel = - new (std::nothrow) kernel::CaffePReluOpenCLKernel(reinterpret_cast(param), inputs, outputs); - if (caffeprelu_kernel == nullptr) { - delete param; - delete input_tensor; - delete output_tensor; - delete weight_tensor; - MS_LOG(ERROR) << "Create caffe prelu kernel error."; - return; - } - - auto ret = caffeprelu_kernel->Init(); - if (ret != RET_OK) { - delete param; - delete input_tensor; - delete output_tensor; - delete weight_tensor; - delete caffeprelu_kernel; - MS_LOG(ERROR) << "caffeprelu_kernel init error."; - return; - } - - MS_LOG(INFO) << "initialize sub_graph"; - std::vector kernels{caffeprelu_kernel}; - auto *sub_graph = new (std::nothrow) kernel::SubGraphOpenCLKernel({input_tensor}, outputs, kernels, kernels, kernels); - if (sub_graph == nullptr) { - delete param; - delete input_tensor; - delete output_tensor; - delete weight_tensor; - delete caffeprelu_kernel; - MS_LOG(ERROR) << "Create sub_graph kernel error."; - return; - } - ret = sub_graph->Init(); - if (ret != RET_OK) { - delete param; - delete input_tensor; - delete output_tensor; - delete weight_tensor; - delete caffeprelu_kernel; - delete sub_graph; - MS_LOG(ERROR) << "sub_graph init error."; - return; - } - MS_LOG(INFO) << "Sub graph begin running!"; - ret = sub_graph->Run(); - if (ret != RET_OK) { - delete input_tensor; - delete output_tensor; - delete weight_tensor; - delete sub_graph; - MS_LOG(ERROR) << "sub_graph run error."; - return; - } - - MS_LOG(INFO) << "CaffePRelu==================output data================"; - printf_tensor_caffeprelu(outputs[0], output_tensor->ElementsC4Num()); - CompareOutCaffePRelu(output_tensor, standard_answer_file); - delete input_tensor; - delete output_tensor; - delete weight_tensor; - delete sub_graph; -} -} // namespace mindspore diff --git a/mindspore/lite/test/ut/src/runtime/kernel/opencl/prelu_tests.cc b/mindspore/lite/test/ut/src/runtime/kernel/opencl/prelu_tests.cc index 91dc9c65596e9a9f0fb72b1951aa4c90a70dfecd..64fc23c43cdf94457c153894978024c925377cd8 100644 --- a/mindspore/lite/test/ut/src/runtime/kernel/opencl/prelu_tests.cc +++ b/mindspore/lite/test/ut/src/runtime/kernel/opencl/prelu_tests.cc @@ -62,26 +62,27 @@ void CompareOutPRelu(lite::tensor::Tensor *output_tensor, const std::string &sta TEST_F(TestPReluOpenCL, PReluFp32_dim4) { std::string in_file = "/data/local/tmp/in_data.bin"; - std::string standard_answer_file = "/data/local/tmp/leaky_relu.bin"; + std::string weight_file = "/data/local/tmp/weight_data.bin"; + std::string standard_answer_file = "/data/local/tmp/caffe_prelu.bin"; MS_LOG(INFO) << "-------------------->> Begin test PRelu!"; auto ocl_runtime = lite::opencl::OpenCLRuntime::GetInstance(); ocl_runtime->Init(); auto allocator = ocl_runtime->GetAllocator(); MS_LOG(INFO) << "Init tensors."; - std::vector input_shape = {1, 4, 3, 8}; + std::vector input_shape = {1, 4, 3, 9}; auto data_type = kNumberTypeFloat32; auto tensor_type = schema::NodeType_ValueNode; auto input_tensor = - new (std::nothrow) lite::tensor::Tensor(data_type, input_shape, schema::Format_NHWC4, tensor_type); + new (std::nothrow) lite::tensor::Tensor(data_type, input_shape, schema::Format_NHWC, tensor_type); if (input_tensor == nullptr) { MS_LOG(ERROR) << "new input_tensor error!"; return; } auto output_tensor = - new (std::nothrow) lite::tensor::Tensor(data_type, input_shape, schema::Format_NHWC4, tensor_type); + new (std::nothrow) lite::tensor::Tensor(data_type, input_shape, schema::Format_NHWC, tensor_type); if (output_tensor == nullptr) { MS_LOG(ERROR) << "new output_tensor error"; delete input_tensor; @@ -89,7 +90,7 @@ TEST_F(TestPReluOpenCL, PReluFp32_dim4) { } auto weight_tensor = - new (std::nothrow) lite::tensor::Tensor(data_type, std::vector{1}, schema::Format_NHWC, tensor_type); + new (std::nothrow) lite::tensor::Tensor(data_type, std::vector{9}, schema::Format_NHWC, tensor_type); if (weight_tensor == nullptr) { MS_LOG(ERROR) << "new weight_tensor error"; delete input_tensor; @@ -105,11 +106,13 @@ TEST_F(TestPReluOpenCL, PReluFp32_dim4) { MS_LOG(INFO) << "initialize input data"; LoadDataPRelu(input_tensor->Data(), input_tensor->Size(), in_file); + LoadDataPRelu(weight_tensor->Data(), weight_tensor->Size(), weight_file); auto weight_data = reinterpret_cast(weight_tensor->Data()); - weight_data[0] = 0.3; + PrintData("Weight data", weight_data, inputs[1]->ElementsNum()); auto *input_data = reinterpret_cast(inputs[0]->Data()); - PrintData("PRelu input data", input_data, inputs[0]->ElementsC4Num()); - + PrintData("PRelu input data", input_data, inputs[0]->ElementsNum()); + std::cout << inputs[0]->ElementsNum() << std::endl; + std::cout << "--------------------------------------------" << std::endl; auto param = new (std::nothrow) PReluParameter(); if (param == nullptr) { MS_LOG(ERROR) << "new PreluParameter error";