!8781 【MSLITE】optimize opencl test code
From: @wangdongxu6 Reviewed-by: Signed-off-by:pull/8781/MERGE
mindspore-ci-bot
5 years ago
committed by
Gitee
commit
a574cdb4e5
@ -1,19 +0,0 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
__constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
|
||||
__kernel void hswish(__read_only image2d_t src_data, __write_only image2d_t dst_data, int4 tensor_shape) {
|
||||
int X = get_global_id(0); // n*h n: default =1
|
||||
int Y = get_global_id(1); // w
|
||||
int Z = get_global_id(2); // c
|
||||
if (X >= tensor_shape.x * tensor_shape.y || Y >= tensor_shape.z || Z >= tensor_shape.w || tensor_shape.y == 0) {
|
||||
return;
|
||||
}
|
||||
int n = X / tensor_shape.y;
|
||||
int h = X % tensor_shape.y;
|
||||
FLT4 temp = READ_IMAGE(src_data, smp_none, (int2)((Y)*tensor_shape.w + Z, (n * tensor_shape.y + h)));
|
||||
FLT4 result = (FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
|
||||
result.x = temp.x <= -3 ? 0 : (temp.x >= 3 ? 1 : temp.x / 6 + 0.5f);
|
||||
result.y = temp.y <= -3 ? 0 : (temp.y >= 3 ? 1 : temp.y / 6 + 0.5f);
|
||||
result.z = temp.z <= -3 ? 0 : (temp.z >= 3 ? 1 : temp.z / 6 + 0.5f);
|
||||
result.w = temp.w <= -3 ? 0 : (temp.w >= 3 ? 1 : temp.w / 6 + 0.5f);
|
||||
WRITE_IMAGE(dst_data, (int2)((Y)*tensor_shape.w + Z, (n * tensor_shape.y + h)), result);
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
@ -1,128 +0,0 @@
|
||||
/**
|
||||
* Copyright 2019 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 "src/runtime/kernel/opencl/kernel/hswish.h"
|
||||
#include <cstring>
|
||||
#include <string>
|
||||
#include <algorithm>
|
||||
#include <set>
|
||||
#include "src/kernel_registry.h"
|
||||
#include "src/runtime/kernel/opencl/utils.h"
|
||||
#include "src/runtime/kernel/opencl/cl/hswish.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_Activation;
|
||||
|
||||
namespace mindspore::kernel {
|
||||
|
||||
int HswishOpenCLKernel::Init() {
|
||||
if (out_tensors_[0]->shape().size() > 4) {
|
||||
MS_LOG(ERROR) << " only support dim <= 4";
|
||||
return RET_ERROR;
|
||||
}
|
||||
|
||||
std::string kernel_name = "hswish";
|
||||
std::set<std::string> build_options;
|
||||
std::string source = hswish_source;
|
||||
std::string program_name = "hswish";
|
||||
ocl_runtime_->LoadSource(program_name, source);
|
||||
ocl_runtime_->BuildKernel(kernel_, program_name, kernel_name, build_options);
|
||||
MS_LOG(DEBUG) << kernel_name << " Init Done!";
|
||||
return RET_OK;
|
||||
}
|
||||
|
||||
void HswishGetWorkGroup(const std::vector<size_t> &global, std::vector<size_t> *local, int max_size) {
|
||||
const int max_divider = 8;
|
||||
const int max_x = 2, max_y = 8;
|
||||
int x = std::min(GetMaxDivisorStrategy1(global[0], max_divider), max_x);
|
||||
int yz = max_size / x;
|
||||
int y = std::min(std::min(GetMaxDivisorStrategy1(global[1], max_divider), yz), max_y);
|
||||
int z = std::min(yz / y, static_cast<int>(UP_DIV(global[2], 2)));
|
||||
|
||||
local->clear();
|
||||
local->push_back(x);
|
||||
local->push_back(y);
|
||||
local->push_back(z);
|
||||
}
|
||||
|
||||
int HswishOpenCLKernel::InferShapeTo4D() {
|
||||
if (in_tensors_[0]->shape().size() <= 4) {
|
||||
if (in_tensors_[0]->shape().size() == 1) {
|
||||
N_ = in_tensors_[0]->shape()[0];
|
||||
} else if (in_tensors_[0]->shape().size() == 2) {
|
||||
N_ = in_tensors_[0]->shape()[0];
|
||||
C_ = in_tensors_[0]->shape()[1];
|
||||
} else if (in_tensors_[0]->shape().size() == 3) {
|
||||
N_ = in_tensors_[0]->shape()[0];
|
||||
W_ = in_tensors_[0]->shape()[1];
|
||||
C_ = in_tensors_[0]->shape()[2];
|
||||
} else {
|
||||
N_ = in_tensors_[0]->shape()[0];
|
||||
H_ = in_tensors_[0]->shape()[1];
|
||||
W_ = in_tensors_[0]->shape()[2];
|
||||
C_ = in_tensors_[0]->shape()[3];
|
||||
}
|
||||
} else {
|
||||
MS_LOG(ERROR) << "Unsupported inputdim: " << in_tensors_[0]->shape().size();
|
||||
return RET_ERROR;
|
||||
}
|
||||
return RET_OK;
|
||||
}
|
||||
|
||||
int HswishOpenCLKernel::Run() {
|
||||
MS_LOG(DEBUG) << this->name() << " Running! ";
|
||||
auto output_shape = out_tensors_[0]->shape();
|
||||
InferShapeTo4D();
|
||||
cl_int4 output_shape_ = {static_cast<cl_int>(N_), static_cast<cl_int>(H_), static_cast<cl_int>(W_),
|
||||
static_cast<cl_int>(UP_DIV(C_, C4NUM))};
|
||||
const std::vector<size_t> &max_global = ocl_runtime_->GetWorkItemSize();
|
||||
std::vector<size_t> local = {1, 1, 1};
|
||||
uint32_t OH = N_ * H_;
|
||||
uint32_t OW = W_;
|
||||
uint32_t OC = UP_DIV(C_, C4NUM);
|
||||
std::vector<size_t> global = {OH, OW, OC};
|
||||
HswishGetWorkGroup(global, &local, max_global[0]);
|
||||
int arg_cn = 0;
|
||||
ocl_runtime_->SetKernelArg(kernel_, arg_cn++, in_tensors_[0]->data_c());
|
||||
ocl_runtime_->SetKernelArg(kernel_, arg_cn++, out_tensors_[0]->data_c());
|
||||
ocl_runtime_->SetKernelArg(kernel_, arg_cn++, output_shape_);
|
||||
ocl_runtime_->RunKernel(kernel_, global, local, nullptr);
|
||||
return RET_OK;
|
||||
}
|
||||
|
||||
kernel::LiteKernel *HswishOpenCLKernelCreator(const std::vector<lite::Tensor *> &inputs,
|
||||
const std::vector<lite::Tensor *> &outputs, OpParameter *opParameter,
|
||||
const lite::InnerContext *ctx, const kernel::KernelKey &desc,
|
||||
const mindspore::lite::PrimitiveC *primitive) {
|
||||
auto *kernel = new (std::nothrow) HswishOpenCLKernel(opParameter, inputs, outputs);
|
||||
if (kernel == nullptr) {
|
||||
MS_LOG(ERROR) << " new HswishOpenCLKernel failed ";
|
||||
free(opParameter);
|
||||
return nullptr;
|
||||
}
|
||||
auto ret = kernel->Init();
|
||||
if (ret != RET_OK) {
|
||||
MS_LOG(ERROR) << " Init kernel failed, name: hswish ";
|
||||
delete kernel;
|
||||
return nullptr;
|
||||
}
|
||||
return kernel;
|
||||
}
|
||||
|
||||
} // namespace mindspore::kernel
|
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue