From 14169f833fb2a8c784fe6544d1b9b03ac32575cd Mon Sep 17 00:00:00 2001 From: liuruilong Date: Mon, 27 Aug 2018 16:40:59 +0800 Subject: [PATCH] update --- .../paddle-mobile/Executor.swift | 2 +- .../paddle-mobile/Operators/BoxcoderOp.swift | 34 ++++++++++++++++--- .../paddle-mobile/Operators/ConcatOp.swift | 2 ++ .../Operators/Kernels/PriorBoxKernel.swift | 10 ++++++ .../Kernels/metal/ReshapeKernel.metal | 22 ++++++------ 5 files changed, 52 insertions(+), 18 deletions(-) diff --git a/metal/paddle-mobile/paddle-mobile/Executor.swift b/metal/paddle-mobile/paddle-mobile/Executor.swift index dcb4a72bd8..33d6650698 100644 --- a/metal/paddle-mobile/paddle-mobile/Executor.swift +++ b/metal/paddle-mobile/paddle-mobile/Executor.swift @@ -14,7 +14,7 @@ import Foundation -let testTo = 93 +let testTo = 94 public class ResultHolder { public let dim: [Int] diff --git a/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift index 994997d491..200b85b490 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift @@ -61,14 +61,38 @@ } func delogOutput() { - let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in - return o - } +// let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in +// return o +// } + + let priorBoxOriginDim = para.priorBox.originDim + let priorBoxArray = para.priorBox.metalTexture.realNHWC(dim: (n: priorBoxOriginDim[0], h: priorBoxOriginDim[1], w: priorBoxOriginDim[2], c: priorBoxOriginDim[3])) + print(" prior box ") + print(priorBoxArray.strideArray()) + + + let priorBoxVarOriginDim = para.priorBoxVar.originDim + let priorBoxVarArray = para.priorBoxVar.metalTexture.realNHWC(dim: (n: priorBoxVarOriginDim[0], h: priorBoxVarOriginDim[1], w: priorBoxVarOriginDim[2], c: priorBoxVarOriginDim[3])) + print(" prior box var ") + print(priorBoxVarArray.strideArray()) + + let targetBoxOriginDim = para.targetBox.originDim + let targetBoxArray = para.targetBox.metalTexture.realNHWC(dim: (n: targetBoxOriginDim[0], h: targetBoxOriginDim[1], w: targetBoxOriginDim[2], c: targetBoxOriginDim[3])) + print(" target box ") + print(targetBoxArray.strideArray()) + + + let originDim = para.output.originDim + + let outputArray = para.output.metalTexture.realNHWC(dim: (n: originDim[0], h: originDim[1], w: originDim[2], c: originDim[3])) print(outputArray.strideArray()) + + +// print(outputArray.strideArray()) //box_coder_0.tmp_0 // writeToLibrary(fileName: "boxcoder_output", array: outputArray) - print(para.output.metalTexture) - print(" write done ") +// print(para.output.metalTexture) +// print(" write done ") } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift index 72711bcff0..986ffbb725 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift @@ -60,6 +60,8 @@ class ConcatOp: Operator, ConcatParam

>, Run func delogOutput() { + + let originDim = para.output.originDim let outputArray = para.output.metalTexture.realNHWC(dim: (n: originDim[0], h: originDim[1], w: originDim[2], c: originDim[3])) print(outputArray.strideArray()) diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift index 09c80f9ab0..77a931a0f5 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift @@ -35,8 +35,18 @@ class PriorBoxKernel: Kernel, Computable{ required init(device: MTLDevice, param: PriorBoxParam

) { super.init(device: device, inFunctionName: "prior_box") param.output.initTexture(device: device, inTranspose: [2, 0, 1, 3]) + + param.outputVariances.initTexture(device: device, inTranspose: [2, 0, 1, 3]) + let n = 1 + let h = param.output.dim[1] + let w = param.output.dim[2] + let c = param.output.dim[3] * param.output.dim[0] + + param.output.dim = Dim.init(inDim: [n, h, w, c]) + param.output.transpose = [0, 1, 2, 3] + let imageWidth = Float32(param.inputImage.originDim[3]) let imageHeight = Float32(param.inputImage.originDim[2]) diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal index 00ba9c1646..533f30156d 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal @@ -77,7 +77,7 @@ kernel void reshape(texture2d_array inTexture [[texture(0)] if (gid.x >= outTexture.get_width() || gid.y >= outTexture.get_height() || gid.z >= outTexture.get_array_size()) return; - + int oxyzn[4] = {int(gid.x), int(gid.y), int(gid.z), 0}, oabcd[4], ixyzn[4], iabcd[4]; ReshapeParam lrp = rp; int oC = lrp.odim[lrp.otrans[3]]; @@ -102,16 +102,14 @@ kernel void reshape(texture2d_array inTexture [[texture(0)] outTexture.write(r, gid.xy, gid.z); } -// -//kernel void reshape_half(texture2d_array inTexture [[texture(0)]], -// texture2d_array outTexture [[texture(1)]], -// uint3 gid [[thread_position_in_grid]]) { -// if (gid.x >= outTexture.get_width() || -// gid.y >= outTexture.get_height() || -// gid.z >= outTexture.get_array_size()) return; -// -// half4 r = inTexture.read(uint2(0, 0), gid.x); -// outTexture.write(r, gid.xy, gid.z); -//} +kernel void reshape_half(texture2d_array inTexture [[texture(0)]], + texture2d_array outTexture [[texture(1)]], + uint3 gid [[thread_position_in_grid]]) { + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height() || + gid.z >= outTexture.get_array_size()) return; + half4 r = inTexture.read(uint2(0, 0), gid.x); + outTexture.write(r, gid.xy, gid.z); +} -- GitLab