提交 8ea39155 编写于 作者: D dolphin8

resize_bilinear

上级 cc2b23d8
develop _release/v2.6.2 gh-pages github/fork/AnBaolei1984/baolei/bitmain github/fork/AnBaolei1984/bitmain github/fork/Cambricon/develop github/fork/DannyIsFunny/Android5 github/fork/DannyIsFunny/Update_MemOpt github/fork/DannyIsFunny/fix_pow github/fork/DannyIsFunny/fix_v26_windows github/fork/GaoWei8/reduce_sum github/fork/GaoWei8/reduce_sum_test_con github/fork/LDOUBLEV/ocr github/fork/Leonardo-Ding/dwh_dev github/fork/MaxwellDing/develop github/fork/MyPandaShaoxiang/fpga_patch github/fork/MyPandaShaoxiang/int8 github/fork/MyPandaShaoxiang/nlp_correct github/fork/MyPandaShaoxiang/opencl_valid github/fork/MyPandaShaoxiang/release/v2.3 github/fork/NHZlX/more_jeston_support github/fork/PaddleLite-EB/merge1.4 github/fork/PaddleLite-EB/new_dev github/fork/Wangzheee/matrix_nms_op github/fork/Xreki/step_rnn/opt_ddim_lite github/fork/cathwong/patch-1 github/fork/cclauss/patch-1 github/fork/chenjiaoAngel/cherry_pic github/fork/chenjiaoAngel/conv_dw_5x5 github/fork/chenjiaoAngel/conv_dw_5x5s2 github/fork/edimetia3d/arm_update_elementwise_op github/fork/edimetia3d/host_deformable_conv github/fork/edimetia3d/matrix_nms_host github/fork/edimetia3d/update_pow_op github/fork/edimetia3d/update_yolo_box github/fork/haozech/develop github/fork/haozech/infershape_chz github/fork/haozech/parl-develop github/fork/jackzhang235/develop github/fork/jameswu2014/develop github/fork/jiansowa/jiansowa/img_nna github/fork/jiweibo/stream_manage github/fork/juncaipeng/add_cast github/fork/lijianshe02/lite-x86 github/fork/qili93/update_sup_model_v26 github/fork/qjing666/develop github/fork/qnqinan/develop github/fork/qnqinan/track-develop github/fork/sangoly/python_compa github/fork/smilejames/develop github/fork/sunsetlh/sunsetlh/xpu_multi_test github/fork/wangqunbaidu/develop github/fork/weihaoji/whj_27 github/fork/weihaoji/xpu_res2net_fusion github/fork/weihaoji/xpu_weihaoji_dev github/fork/xiebaiyuan/fix_leak_opencl github/fork/xiebaiyuan/opencl_depthwised1 github/fork/xiebaiyuan/opencl_softmax github/fork/yanghongtian/yanghongtian/add_ascend310_target_place github/fork/yiicy/computelib github/fork/yongqiangma/bm_card github/fork/yongqiangma/calib github/fork/yongqiangma/copytocpu github/fork/yongqiangma/gpu github/fork/yongqiangma/pass github/fork/yongqiangma/pool github/fork/yongqiangma/priorbox github/fork/yongqiangma/shape github/fork/yongqiangma/split_c github/fork/yongqiangma/trans github/fork/yongqiangma/trans2 github/fork/yongqiangma/workspace github/fork/ysh329/add-cl-kernel-member-for-opencl github/fork/ysh329/add-get-output github/fork/ysh329/cherry-pick-precision-profiler-enhance github/fork/ysh329/fix-opencl-concat github/fork/ysh329/support-int64-copy-from-to-cpu github/fork/zhaoyang-star/enable_prifile_in_tiny_publish github/fork/zhaoyang-star/fix_openc_demo github/fork/zhaoyang-star/patch-1 github/fork/zhupengyang/opt release/v2.0.0 release/v2.0.0-beta1 release/v2.0.0-beta2 release/v2.0.0-rc release/v2.1.0 release/v2.2.0 release/v2.3 release/v2.6 release/v2.6.0 release/v2.7 revert-4368-hongming/test_v26 2.0.0-beta 1.5.0 1.1.1 1.1.0 v2.7-beta v2.6.3-beta2 v2.6.3-beta1 v2.6.2 v2.6.1 v2.6.0 v2.3.0 v2.2.0 v2.1.0 v2.0.0 v2.0.0-rc v2.0.0-beta1 v2.0.0-beta1-prerel v1.0 release/1.4
8 合并请求!3489pull code,!3210[Opencl] fix opencl bug,!3154[arm]resize nnv12 bug,!3074[opencl]add grid_sampler op,!1241fix #1239 release 1.1.0,!1011fix #1010 merge new version metal,!950texture,!854Metal
......@@ -60,7 +60,8 @@ class OpCreator<P: PrecisionType> {
gTransposeType : TransposeOp<P>.creat,
gPriorBoxType : PriorBoxOp<P>.creat,
gPreluType : PreluOp<P>.creat,
gConv2dTransposeType : ConvTransposeOp<P>.creat]
gConv2dTransposeType : ConvTransposeOp<P>.creat,
gResizeBilinearType : ResizeBilinearOp<P>.creat]
private init(){}
}
......@@ -139,6 +139,7 @@ let gConvBnReluType = "conv_bn_relu"
let gDwConvBnReluType = "depth_conv_bn_relu"
let gPreluType = "prelu"
let gConv2dTransposeType = "conv2d_transpose"
let gResizeBilinearType = "resize_bilinear"
let opInfos = [gConvType : (inputs: ["Input"], outputs: ["Output"]),
......@@ -161,5 +162,6 @@ let opInfos = [gConvType : (inputs: ["Input"], outputs: ["Out
gMulticlassNMSType : (inputs: ["BBoxes", "Scores"], outputs: ["Out"]),
gPriorBoxType : (inputs: ["Input", "Image"], outputs: ["Boxes", "Variances"]),
gPreluType : (inputs: ["X"], outputs: ["Out"]),
gConv2dTransposeType : (inputs: ["Input"], outputs: ["Output"])
gConv2dTransposeType : (inputs: ["Input"], outputs: ["Output"]),
gResizeBilinearType : (inputs: ["X"], outputs: ["Out"])
]
/* 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. */
import Foundation
struct ResizeBilinearMetalParam {
var ratio_h: Float32
var ratio_w: Float32
}
class ResizeBilinearKernel<P: PrecisionType>: Kernel, Computable{
func compute(commandBuffer: MTLCommandBuffer, param: ResizeBilinearParam<P>) throws {
guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
throw PaddleMobileError.predictError(message: " encode is nil")
}
encoder.setTexture(param.input.metalTexture, index: 0)
encoder.setTexture(param.output.metalTexture, index: 1)
let ratio_h: Float32 = Float32(param.input.tensorDim.dims[2]) / Float32(param.output.tensorDim.dims[2])
let ratio_w: Float32 = Float32(param.input.tensorDim.dims[3]) / Float32(param.output.tensorDim.dims[3])
var p = ResizeBilinearMetalParam.init(ratio_h: ratio_h, ratio_w: ratio_w)
encoder.setBytes(&p, length: MemoryLayout<ConcatMetalParam>.size, index: 0)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
}
required init(device: MTLDevice, param: ResizeBilinearParam<P>) {
param.output.initTexture(device: device, inTranspose: param.input.transpose, computePrecision: computePrecision)
if computePrecision == .Float32 {
super.init(device: device, inFunctionName: "resize_bilinear")
} else if computePrecision == .Float16 {
super.init(device: device, inFunctionName: "resize_bilinear_half")
} else {
fatalError()
}
}
}
/* 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. */
#include <metal_stdlib>
using namespace metal;
struct resize_bilinear_param {
// int32_t out_h;
// int32_t out_w;
float ratio_h;
float ratio_w;
};
kernel void resize_bilinear(texture2d_array<float, access::read> input [[texture(0)]],
texture2d_array<float, access::write> output [[texture(2)]],
constant resize_bilinear_param & pm [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
float4 r;
if ((input.get_width() == output.get_width()) && (input.get_height() == output.get_height())) {
r = input.read(gid.xy, gid.z)
} else {
float w = gid.x * pm.ratio_w;
float h = gid.y * pm.ratio_h;
uint w0 = w, h0 = h;
uint w1 = w0 + 1, h1 = h0 + 1;
float w1lambda = w - w0, h1lambda = h - h0;
float w2lambda = 1.0 - w1lambda, h2lambda = 1.0 - h1lambda;
if (w1 >= input.get_width()) w1 = w0;
if (h1 >= input.get_height()) h1 = h0;
float4 r0 = input.read(uint2(w0, h0), gid.z);
float4 r1 = input.read(uint2(w1, h0), gid.z);
float4 r2 = input.read(uint2(w0, h1), gid.z);
float4 r3 = input.read(uint2(w1, h1), gid.z);
r = h2lambda * (w2lambda * r0 + w1lambda * r1) + h1lambda * (w2lambda * r3 + w1lambda * r4);
}
output.write(r, gid.xy, gid.z);
}
kernel void resize_bilinear_half(texture2d_array<half, access::read> input [[texture(0)]],
texture2d_array<half, access::write> output [[texture(2)]],
constant resize_bilinear_param & pm [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
half4 r;
if ((input.get_width() == output.get_width()) && (input.get_height() == output.get_height())) {
r = input.read(gid.xy, gid.z)
} else {
half w = gid.x * pm.ratio_w;
half h = gid.y * pm.ratio_h;
uint w0 = w, h0 = h;
uint w1 = w0 + 1, h1 = h0 + 1;
half w1lambda = w - w0, h1lambda = h - h0;
half w2lambda = 1.0 - w1lambda, h2lambda = 1.0 - h1lambda;
if (w1 >= input.get_width()) w1 = w0;
if (h1 >= input.get_height()) h1 = h0;
half4 r0 = input.read(uint2(w0, h0), gid.z);
half4 r1 = input.read(uint2(w1, h0), gid.z);
half4 r2 = input.read(uint2(w0, h1), gid.z);
half4 r3 = input.read(uint2(w1, h1), gid.z);
r = h2lambda * (w2lambda * r0 + w1lambda * r1) + h1lambda * (w2lambda * r3 + w1lambda * r4);
}
output.write(r, gid.xy, gid.z);
output.write(r, gid.xy, gid.z);
}
///* 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. */
import Foundation
class ResizeBilinearParam<P: PrecisionType>: OpParam {
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
input = try ResizeBilinearParam.inputX(inputs: opDesc.inputs, from: inScope)
if (input.transpose != [0, 2, 3, 1]) || (input.tensorDim.cout() != 4) {
fatalError()
}
output = try ResizeBilinearParam.outputOut(outputs: opDesc.outputs, from: inScope)
out_h = try ResizeBilinearParam.getAttr(key: "out_h", attrs: opDesc.attrs)
out_w = try ResizeBilinearParam.getAttr(key: "out_w", attrs: opDesc.attrs)
} catch let error {
throw error
}
}
let input: Texture<P>
var output: Texture<P>
let out_h: Int32
let out_w: Int32
}
class ResizeBilinearOp<P: PrecisionType>: Operator<ResizeBilinearKernel<P>, ResizeBilinearParam<P>>, Runable, Creator, InferShaperable{
typealias OpType = ResizeBilinearOp<P>
func inferShape() {
// para.output.dim = para.input.dim
}
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do {
try kernel.compute(commandBuffer: buffer, param: para)
} catch let error {
throw error
}
}
func delogOutput() {
print(" \(type) output: ")
}
}
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册