diff --git a/mindspore/lite/src/runtime/kernel/opencl/cl/activation.cl b/mindspore/lite/src/runtime/kernel/opencl/cl/activation.cl index 7698a0228acc8b33180782b7fb0899a15a44893b..f745568c2d40b838b10ffd30544b9c9ed9acdb63 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/cl/activation.cl +++ b/mindspore/lite/src/runtime/kernel/opencl/cl/activation.cl @@ -5,33 +5,18 @@ #define MIN(X, Y) (X < Y ? X : Y) __constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; -__kernel void LeakyRelu_NHWC4(__read_only image2d_t input, __write_only image2d_t output, const int4 img_shape, - __global FLT4 *alpha, const int4 input_shape) { +__kernel void LeakyRelu(__read_only image2d_t input, __write_only image2d_t output, const int4 img_shape, + const float alpha) { int Y = get_global_id(0); // H int X = get_global_id(1); // W C4 if (X >= img_shape.z || Y >= img_shape.y) return; - int C = X % UP_DIV(input_shape.w, SLICES); FLT4 in_c4 = READ_IMAGE(input, smp_zero, (int2)(X, Y)); FLT4 tmp; - tmp.x = in_c4.x > 0.0f ? in_c4.x : in_c4.x * alpha[C].x; - tmp.y = in_c4.y > 0.0f ? in_c4.y : in_c4.y * alpha[C].y; - tmp.z = in_c4.z > 0.0f ? in_c4.z : in_c4.z * alpha[C].z; - tmp.w = in_c4.w > 0.0f ? in_c4.w : in_c4.w * alpha[C].w; - WRITE_IMAGE(output, (int2)(X, Y), tmp); -} - -__kernel void LeakyRelu_NC4HW4(__read_only image2d_t input, __write_only image2d_t output, const int4 img_shape, - __global FLT4 *alpha, const int4 input_shape) { - int Y = get_global_id(0); // C4 H - int X = get_global_id(1); // W - if (X >= img_shape.z || Y >= img_shape.y) return; - int C = Y / input_shape.y; - FLT4 in_c4 = READ_IMAGE(input, smp_zero, (int2)(X, Y)); - FLT4 tmp; - tmp.x = in_c4.x > 0.0f ? in_c4.x : in_c4.x * alpha[C].x; - tmp.y = in_c4.y > 0.0f ? in_c4.y : in_c4.y * alpha[C].y; - tmp.z = in_c4.z > 0.0f ? in_c4.z : in_c4.z * alpha[C].z; - tmp.w = in_c4.w > 0.0f ? in_c4.w : in_c4.w * alpha[C].w; + FLT alpha_f = TO_FLT(alpha); + tmp.x = in_c4.x > 0.0f ? in_c4.x : in_c4.x * alpha_f; + tmp.y = in_c4.y > 0.0f ? in_c4.y : in_c4.y * alpha_f; + tmp.z = in_c4.z > 0.0f ? in_c4.z : in_c4.z * alpha_f; + tmp.w = in_c4.w > 0.0f ? in_c4.w : in_c4.w * alpha_f; WRITE_IMAGE(output, (int2)(X, Y), tmp); } diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/activation.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/activation.cc index 4030148870f37d572776838e96064df24fb09bdd..c4dfae3c462d26f64abbdee1e178fe09fea1ff94 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/activation.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/activation.cc @@ -39,46 +39,7 @@ using mindspore::schema::PrimitiveType_Activation; namespace mindspore::kernel { -void ActivationOpenClKernel::InitBuffer() { - auto allocator = lite::opencl::OpenCLRuntime::GetInstance()->GetAllocator(); - int elem_num = UP_ROUND(nhwc_shape_[3], C4NUM); - alpha_buff_ = allocator->Malloc(elem_num * fp_size); - alpha_buff_ = allocator->MapBuffer(alpha_buff_, CL_MAP_WRITE, nullptr, true); - memset(alpha_buff_, 0x00, elem_num * fp_size); - if (in_tensors_.size() == 1) { - if (enable_fp16_) { - uint16_t alpha_fp16 = Float32ToShort(alpha_); - auto alpha_buff_fp16 = reinterpret_cast(alpha_buff_); - for (int i = 0; i < nhwc_shape_[3]; i++) { - alpha_buff_fp16[i] = alpha_fp16; - } - } else { - auto alpha_buff_fp16 = reinterpret_cast(alpha_buff_); - for (int i = 0; i < nhwc_shape_[3]; i++) { - alpha_buff_fp16[i] = alpha_; - } - } - } else { - if (enable_fp16_) { - if (in_tensors_[1]->data_type() == kNumberTypeFloat32) { - auto alpha_buff_fp16 = reinterpret_cast(alpha_buff_); - for (int i = 0; i < nhwc_shape_[3]; i++) { - alpha_buff_fp16[i] = Float32ToShort(reinterpret_cast(in_tensors_[0]->Data())[i]); - } - } else { - memcpy(alpha_buff_, in_tensors_[0]->Data(), nhwc_shape_[3] * fp_size); - } - } else { - if (in_tensors_[1]->data_type() == kNumberTypeFloat16) { - MS_LOG(WARNING) << "fp16 model run in fp32 mode not support."; - memcpy(alpha_buff_, in_tensors_[0]->Data(), nhwc_shape_[3] * fp_size); - } else { - memcpy(alpha_buff_, in_tensors_[0]->Data(), nhwc_shape_[3] * fp_size); - } - } - } - allocator->UnmapBuffer(alpha_buff_); -} +void ActivationOpenClKernel::InitBuffer() {} int ActivationOpenClKernel::Init() { in_size_ = in_tensors_[0]->shape().size(); @@ -102,9 +63,6 @@ int ActivationOpenClKernel::Init() { MS_LOG(ERROR) << "Activate fun only support dim=4 or 2, but your dim=" << in_size_; return RET_ERROR; } - if (type_ == ActivationType_LEAKY_RELU) { - InitBuffer(); - } std::map> Program_Kernel{ {ActivationType_LEAKY_RELU, std::vector{"LEAKY_RELU", "LeakyRelu"}}, {ActivationType_RELU, std::vector{"RELU", "Relu"}}, @@ -119,9 +77,6 @@ int ActivationOpenClKernel::Init() { std::set build_options; ocl_runtime->LoadSource(Program_Kernel[type_][0], source); std::string kernel_name = Program_Kernel[type_][1]; - if (type_ == ActivationType_LEAKY_RELU) { - kernel_name += "_" + std::string(EnumNameFormat(op_format_)); - } ocl_runtime->BuildKernel(kernel_, Program_Kernel[type_][0], kernel_name, build_options); in_ori_format_ = in_tensors_[0]->GetFormat(); out_ori_format_ = out_tensors_[0]->GetFormat(); @@ -140,10 +95,7 @@ int ActivationOpenClKernel::Run() { ocl_runtime->SetKernelArg(kernel_, arg_idx++, out_tensors_[0]->Data()); ocl_runtime->SetKernelArg(kernel_, arg_idx++, img2d_shape); if (type_ == ActivationType_LEAKY_RELU) { - ocl_runtime->SetKernelArg(kernel_, arg_idx++, alpha_buff_, lite::opencl::MemType::BUF); - cl_int4 input_shape = {static_cast(nhwc_shape_[0]), static_cast(nhwc_shape_[1]), - static_cast(nhwc_shape_[2]), static_cast(nhwc_shape_[3])}; - ocl_runtime->SetKernelArg(kernel_, arg_idx++, input_shape); + ocl_runtime->SetKernelArg(kernel_, arg_idx++, alpha_); } std::vector local = {}; std::vector global = {static_cast(img2d_shape.s[1]), static_cast(img2d_shape.s[2])}; diff --git a/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc b/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc index e5a7d35a8d84d39604332683608f4ba29a06d470..fb008eddd0fd2efac748da31b9a1ce6d64c2edc1 100644 --- a/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc +++ b/mindspore/lite/src/runtime/kernel/opencl/kernel/prelu.cc @@ -22,6 +22,7 @@ #include "src/kernel_registry.h" #include "include/errorcode.h" +#include "nnacl/fp32/common_func.h" #include "src/runtime/kernel/opencl/kernel/prelu.h" #include "src/runtime/opencl/opencl_runtime.h" #include "src/runtime/kernel/opencl/cl/prelu.cl.inc" @@ -35,18 +36,38 @@ 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); auto allocator = lite::opencl::OpenCLRuntime::GetInstance()->GetAllocator(); + int elem_num = in_tensors_[0]->shape().size() == 2 ? in_tensors_[0]->shape()[1] : in_tensors_[0]->shape()[3]; + int elem_num_c4 = UP_DIV(elem_num, C4NUM); size_t img_dtype = CL_FLOAT; if (enable_fp16_) { img_dtype = CL_HALF_FLOAT; } - std::vector img_size{size_t(div_ci), 1, img_dtype}; - PReluWeight_ = allocator->Malloc(div_ci * C4NUM * fp_size, img_size); + std::vector img_size{size_t(elem_num_c4), 1, img_dtype}; + PReluWeight_ = allocator->Malloc(elem_num_c4 * C4NUM * fp_size, img_size); PReluWeight_ = allocator->MapBuffer(PReluWeight_, CL_MAP_WRITE, nullptr, true); - memset(PReluWeight_, 0x00, div_ci * C4NUM * fp_size); - memcpy(PReluWeight_, in_tensors_[1]->Data(), C * fp_size); + memset(PReluWeight_, 0x00, elem_num_c4 * C4NUM * fp_size); + if (enable_fp16_) { + if (in_tensors_[1]->data_type() == kNumberTypeFloat32) { + auto PReluWeight_fp16 = reinterpret_cast(PReluWeight_); + auto in_tensor_data_fp32 = reinterpret_cast(in_tensors_[1]->Data()); + for (int i = 0; i < elem_num; i++) { + PReluWeight_fp16[i] = Float32ToShort(in_tensor_data_fp32[i]); + } + } else { + memcpy(PReluWeight_, in_tensors_[1]->Data(), elem_num * fp_size); + } + } else { + if (in_tensors_[1]->data_type() == kNumberTypeFloat16) { + auto PReluWeight_fp32 = reinterpret_cast(PReluWeight_); + auto in_tensor_data_fp16 = reinterpret_cast(in_tensors_[1]->Data()); + for (int i = 0; i < elem_num; i++) { + PReluWeight_fp32[i] = ShortToFloat32(in_tensor_data_fp16[i]); + } + } else { + memcpy(PReluWeight_, in_tensors_[1]->Data(), elem_num * fp_size); + } + } allocator->UnmapBuffer(PReluWeight_); } diff --git a/mindspore/lite/test/ut/src/runtime/kernel/opencl/activation_tests.cc b/mindspore/lite/test/ut/src/runtime/kernel/opencl/activation_tests.cc index e95c7fa8388c65b1f0ffd6851a2b8597d6038742..19422bafe11cea4cab919a1112a47d7e765d6700 100644 --- a/mindspore/lite/test/ut/src/runtime/kernel/opencl/activation_tests.cc +++ b/mindspore/lite/test/ut/src/runtime/kernel/opencl/activation_tests.cc @@ -432,7 +432,7 @@ TEST_F(TestActivationOpenCL, LeakyReluFp_dim4) { std::vector input_shape = {1, 9}; // need modify auto tensor_type = schema::NodeType_ValueNode; schema::Format format = schema::Format_NC; // need modify - schema::Format op_format = schema::Format_NC4; // need modify + schema::Format op_format = schema::Format_NHWC4; // need modify auto *input_tensor = new (std::nothrow) lite::tensor::Tensor(data_type, input_shape, format, tensor_type); if (input_tensor == nullptr) { MS_LOG(ERROR) << "new input tensor error!";