diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj b/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj index f3ab9fc66a072cd5b0bbba56ae99258f04be3612..6f3db873591d1d4ca05fc1ea563d01f44c779d7e 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj +++ b/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj @@ -14,9 +14,11 @@ FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8520E11C550081E9F8 /* Main.storyboard */; }; FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8820E11C560081E9F8 /* Assets.xcassets */; }; FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */; }; - FC3602C82108580600FACB58 /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3602C72108580600FACB58 /* MetalHelper.swift */; }; FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */ = {isa = PBXBuildFile; fileRef = FC918190211DBC3500B6F354 /* paddle-mobile.png */; }; FC918193211DC70500B6F354 /* iphone.JPG in Resources */ = {isa = PBXBuildFile; fileRef = FC918192211DC70500B6F354 /* iphone.JPG */; }; + FCBCCC522122EEDC00D94F7E /* ssd_hand_params in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC502122EEDC00D94F7E /* ssd_hand_params */; }; + FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC512122EEDC00D94F7E /* ssd_hand_model */; }; + FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC542122EF5400D94F7E /* MetalHelper.swift */; }; FCD04E6320F3146B0007374F /* params in Resources */ = {isa = PBXBuildFile; fileRef = FCD04E6120F3146A0007374F /* params */; }; FCD04E6420F3146B0007374F /* model in Resources */ = {isa = PBXBuildFile; fileRef = FCD04E6220F3146A0007374F /* model */; }; FCDFD3FB211D72C3005AB38B /* ModelHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCDFD3FA211D72C3005AB38B /* ModelHelper.swift */; }; @@ -52,9 +54,11 @@ FC039B8820E11C560081E9F8 /* Assets.xcassets */ = {isa = PBXFileReference; lastKnownFileType = folder.assetcatalog; path = Assets.xcassets; sourceTree = ""; }; FC039B8B20E11C560081E9F8 /* Base */ = {isa = PBXFileReference; lastKnownFileType = file.storyboard; name = Base; path = Base.lproj/LaunchScreen.storyboard; sourceTree = ""; }; FC039B8D20E11C560081E9F8 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; sourceTree = ""; }; - FC3602C72108580600FACB58 /* MetalHelper.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; name = MetalHelper.swift; path = "../../paddle-mobile-unit-test/paddle-mobile-unit-test/MetalHelper.swift"; sourceTree = ""; }; FC918190211DBC3500B6F354 /* paddle-mobile.png */ = {isa = PBXFileReference; lastKnownFileType = image.png; path = "paddle-mobile.png"; sourceTree = ""; }; FC918192211DC70500B6F354 /* iphone.JPG */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = iphone.JPG; sourceTree = ""; }; + FCBCCC502122EEDC00D94F7E /* ssd_hand_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_params; sourceTree = ""; }; + FCBCCC512122EEDC00D94F7E /* ssd_hand_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_model; sourceTree = ""; }; + FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = ""; }; FCD04E6120F3146A0007374F /* params */ = {isa = PBXFileReference; lastKnownFileType = file; path = params; sourceTree = ""; }; FCD04E6220F3146A0007374F /* model */ = {isa = PBXFileReference; lastKnownFileType = file; path = model; sourceTree = ""; }; FCDFD3FA211D72C3005AB38B /* ModelHelper.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ModelHelper.swift; sourceTree = ""; }; @@ -124,7 +128,7 @@ FC039B8820E11C560081E9F8 /* Assets.xcassets */, FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */, FC039B8D20E11C560081E9F8 /* Info.plist */, - FC3602C72108580600FACB58 /* MetalHelper.swift */, + FCBCCC542122EF5400D94F7E /* MetalHelper.swift */, FCDFD3FA211D72C3005AB38B /* ModelHelper.swift */, ); path = "paddle-mobile-demo"; @@ -145,12 +149,22 @@ FC0E2C2020EDC03B009C1FAC /* models */ = { isa = PBXGroup; children = ( + FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */, FCD04E6020F3146A0007374F /* mobilenet */, ); name = models; path = ../../models; sourceTree = ""; }; + FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */ = { + isa = PBXGroup; + children = ( + FCBCCC502122EEDC00D94F7E /* ssd_hand_params */, + FCBCCC512122EEDC00D94F7E /* ssd_hand_model */, + ); + path = mobilenet_ssd_hand; + sourceTree = ""; + }; FCD04E6020F3146A0007374F /* mobilenet */ = { isa = PBXGroup; children = ( @@ -225,11 +239,13 @@ FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */, FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */, FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */, + FCBCCC522122EEDC00D94F7E /* ssd_hand_params in Resources */, FCEEE7D4210627A000444BEC /* banana.jpeg in Resources */, FC918193211DC70500B6F354 /* iphone.JPG in Resources */, FCDFD41B211D91C7005AB38B /* synset.txt in Resources */, FCD04E6420F3146B0007374F /* model in Resources */, FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */, + FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */, ); runOnlyForDeploymentPostprocessing = 0; }; @@ -282,8 +298,8 @@ FC039B8420E11C550081E9F8 /* ViewController.swift in Sources */, FCDFD3FB211D72C3005AB38B /* ModelHelper.swift in Sources */, FC013928210204A3008100E3 /* PreProcessKernel.metal in Sources */, + FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */, FC039B8220E11C550081E9F8 /* AppDelegate.swift in Sources */, - FC3602C82108580600FACB58 /* MetalHelper.swift in Sources */, ); runOnlyForDeploymentPostprocessing = 0; }; diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift index 7e1f66855e45453eee9fdbe034a309aee44ff960..53a15685fe3a875bc24fda91fad6f5d8cf6e7405 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift @@ -12,23 +12,17 @@ import Foundation import paddle_mobile import MetalPerformanceShaders -class PreProccess: CusomKernel { - init(device: MTLDevice) { - let s = CusomKernel.Shape.init(inWidth: 224, inHeight: 224, inChannel: 3) - super.init(device: device, inFunctionName: "preprocess", outputDim: s, usePaddleMobileLib: false) - } -} - -let modelHelperMap: [SupportModel : ModelHelper] = [.mobilenet : MobileNetHelper.init()] +let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()] enum SupportModel: String{ case mobilenet = "mobilenet" + case mobilenet_ssd = "mobilenetssd" static func supportedModels() -> [SupportModel] { - return [.mobilenet] + return [.mobilenet, .mobilenet_ssd] } } -protocol ModelHelper { +protocol Net { var dim: [Int] { get } var modelPath: String { get } var paramPath: String { get } @@ -38,7 +32,7 @@ protocol ModelHelper { func resultStr(res: [Float]) -> String } -extension ModelHelper { +extension Net { func getTexture(image: CGImage, getTexture: @escaping (MTLTexture) -> Void) { let texture = try? MetalHelper.shared.textureLoader.newTexture(cgImage: image, options: [:]) ?! " texture loader error" MetalHelper.scaleTexture(queue: MetalHelper.shared.queue, input: texture!, size: (224, 224)) { (resTexture) in @@ -47,7 +41,15 @@ extension ModelHelper { } } -struct MobileNetHelper: ModelHelper{ +struct MobileNet: Net{ + + class MobilenetPreProccess: CusomKernel { + init(device: MTLDevice) { + let s = CusomKernel.Shape.init(inWidth: 224, inHeight: 224, inChannel: 3) + super.init(device: device, inFunctionName: "preprocess", outputDim: s, usePaddleMobileLib: false) + } + } + class PreWords { var contents: [String] = [] init(fileName: String, type: String = "txt", inBundle: Bundle = Bundle.main) { @@ -84,6 +86,33 @@ struct MobileNetHelper: ModelHelper{ modelPath = Bundle.main.path(forResource: "model", ofType: nil) ?! "model null" paramPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null" modelDir = "" - preprocessKernel = PreProccess.init(device: MetalHelper.shared.device) + preprocessKernel = MobilenetPreProccess.init(device: MetalHelper.shared.device) + } +} + +struct MobileNet_ssd_hand: Net{ + class MobilenetssdPreProccess: CusomKernel { + init(device: MTLDevice) { + let s = CusomKernel.Shape.init(inWidth: 300, inHeight: 300, inChannel: 3) + super.init(device: device, inFunctionName: "mobilenet_ssd_preprocess", outputDim: s, usePaddleMobileLib: false) + } + } + + func resultStr(res: [Float]) -> String { + fatalError() + } + + var preprocessKernel: CusomKernel + let dim = [1, 300, 300, 3] + let modelPath: String + let paramPath: String + let modelDir: String + + init() { + modelPath = Bundle.main.path(forResource: "ssd_hand_model", ofType: nil) ?! "model null" + paramPath = Bundle.main.path(forResource: "ssd_hand_params", ofType: nil) ?! "para null" + modelDir = "" + preprocessKernel = MobilenetssdPreProccess.init(device: MetalHelper.shared.device) } } + diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/PreProcessKernel.metal b/metal/paddle-mobile-demo/paddle-mobile-demo/PreProcessKernel.metal index f359ab39ac5fbc18febfb6f0da367e72b61b959c..ab43f8d9233c581cd6957befb2decafc557f1264 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/PreProcessKernel.metal +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/PreProcessKernel.metal @@ -39,6 +39,34 @@ kernel void preprocess_half( } +kernel void mobilenet_ssd_preprocess( + texture2d inTexture [[texture(0)]], + texture2d outTexture [[texture(1)]], + uint2 gid [[thread_position_in_grid]]) +{ + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height()) { + return; + } + const auto means = float4(123.68f, 116.78f, 103.94f, 0.0f); + const float4 inColor = (inTexture.read(gid) * 255.0 - means) * 0.017; + outTexture.write(float4(inColor.z, inColor.y, inColor.x, 0.0f), gid); +} + +kernel void mobilenet_ssd_preprocess_half( + texture2d inTexture [[texture(0)]], + texture2d outTexture [[texture(1)]], + uint2 gid [[thread_position_in_grid]]) +{ + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height()) { + return; + } + const auto means = half4(123.68f, 116.78f, 103.94f, 0.0f); + const half4 inColor = (inTexture.read(gid) * 255.0 - means) * 0.017; + outTexture.write(half4(inColor.z, inColor.y, inColor.x, 0.0f), gid); +} + diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift index 30fdaf078556bdc4546aec4f27e153f469d9e5ac..241b50e2efdb664e59bf83c43e9a0e58cc68aece 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift @@ -30,7 +30,7 @@ class ViewController: UIViewController { var executor: Executor? var modelType: SupportModel = .mobilenet var toPredictTexture: MTLTexture? - var modelHelper: ModelHelper { + var modelHelper: Net { return modelHelperMap[modelType] ?! " has no this type " } var threadNum = 1 @@ -76,7 +76,7 @@ class ViewController: UIViewController { } do { - let max = 100 + let max = 10 var startDate = Date.init() for i in 0.. { [gConvType : ConvOp

.creat, gBatchNormType : BatchNormOp

.creat, gReluType : ReluOp

.creat, - gElementwiseAdd : ElementwiseAddOp

.creat, + gElementwiseAddType : ElementwiseAddOp

.creat, gFeedType : FeedOp

.creat, gFetchType : FetchOp

.creat, gConvAddBatchNormReluType : ConvAddBatchNormReluOp

.creat, gPooType : PoolOp

.creat, gSoftmaxType : SoftmaxOp

.creat, gReshapeType : ReshapeOp

.creat, - gConvAddType : ConvAddOp

.creat] + gConvAddType : ConvAddOp

.creat, + gDepthConvType : DepthConvOp

.creat, + gConcatType : ConcatOp

.creat, + gBoxcoderType : BoxcoderOp

.creat, + gConvBnReluType : ConvBNReluOp

.creat, + gDwConvBnReluType : DwConvBNReluOp

.creat, + gMulticlassNMSType : MulticlassNMSOp

.creat] private init(){} } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift b/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift index bc95f84d8ae98cb8e4e7151f0cf69a574699dc80..952004a00d4fcac70c45164893a9b35f8d177b9a 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift @@ -118,22 +118,37 @@ let gFeedType = "feed" let gConvType = "conv2d" let gBatchNormType = "batch_norm" let gReluType = "relu" -let gElementwiseAdd = "elementwise_add" +let gElementwiseAddType = "elementwise_add" let gConvAddBatchNormReluType = "conv_add_batchnorm_relu" let gPooType = "pool2d" let gSoftmaxType = "softmax" let gReshapeType = "reshape" let gConvAddType = "conv_add" - +let gDepthConvType = "depthwise_conv2d" +let gPriorBoxType = "prior_box" +let gTransposeType = "transpose" +let gConcatType = "concat" +let gBoxcoderType = "box_coder" +let gMulticlassNMSType = "multiclass_nms" +let gConvBnReluType = "conv_bn_relu" +let gDwConvBnReluType = "depth_conv_bn_relu" let opInfos = [gConvType : (inputs: ["Input"], outputs: ["Output"]), gBatchNormType : (inputs: ["X"], outputs: ["Y"]), gReluType : (inputs: ["X"], outputs: ["Out"]), - gElementwiseAdd : (inputs: ["X"], outputs: ["Out"]), + gElementwiseAddType : (inputs: ["X"], outputs: ["Out"]), gFeedType : (inputs: ["X"], outputs: ["Out"]), gFetchType : (inputs: ["X"], outputs: ["Out"]), gConvAddBatchNormReluType : (inputs: ["Input"], outputs: ["Out"]), gPooType : (inputs: ["X"], outputs: ["Out"]), gSoftmaxType : (inputs: ["X"], outputs: ["Out"]), gReshapeType : (inputs: ["X"], outputs: ["Out"]), - gConvAddType : (inputs: ["Input"], outputs: ["Out"])] + gConvAddType : (inputs: ["Input"], outputs: ["Out"]), + gDepthConvType : (inputs: ["Input"], outputs: ["Output"]), + gConcatType : (inputs: ["X"], outputs: ["Out"]), + gBoxcoderType : (inputs: ["PriorBox", "PriorBoxVar", "TargetBox"], outputs: ["OutputBox"]), + gTransposeType : (inputs: ["X"], outputs: ["Out"]), + gConvBnReluType : (inputs: ["Input"], outputs: ["Out"]), + gDwConvBnReluType : (inputs: ["Input"], outputs: ["Out"]), + gMulticlassNMSType : (inputs: ["BBoxes", "Scores"], outputs: ["Out"]), + gPriorBoxType : (inputs: ["Input", "Image"], outputs: ["Boxes", "Variances"])] diff --git a/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift new file mode 100644 index 0000000000000000000000000000000000000000..8eeae8867ec6f40b0dadaf5ba96aa847516f368f --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift @@ -0,0 +1,47 @@ +///* 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 BoxcoderParam: OpParam { + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + fatalError() + } catch let error { + throw error + } + } + let input: Texture

+ var output: Texture

+} + +class BoxcoderOp: Operator, BoxcoderParam

>, Runable, Creator, InferShaperable{ + + func inferShape() { + para.output.dim = para.input.dim + } + + typealias OpType = BoxcoderOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } + } +} + + + diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift new file mode 100644 index 0000000000000000000000000000000000000000..a85fec5880738791abc3551f5b712a736ba1f46d --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift @@ -0,0 +1,47 @@ +///* 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 ConcatParam: OpParam { + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + fatalError() + } catch let error { + throw error + } + } + let input: Texture

+ var output: Texture

+} + +class ConcatOp: Operator, ConcatParam

>, Runable, Creator, InferShaperable{ + + func inferShape() { + para.output.dim = para.input.dim + } + + typealias OpType = ConcatOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } + } +} + + + diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift index f24e25b054f02c7b8f12015697fd61e9a2005ef8..4814a63f3675fb4f246950dd43b82387cf42cc60 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift @@ -92,7 +92,7 @@ class ConvAddBatchNormReluOp: Operator Node { let beginNode = Node.init(inType: gConvType) _ = beginNode - --> Node.init(inType: gElementwiseAdd) + --> Node.init(inType: gElementwiseAddType) --> Node.init(inType: gBatchNormType) --> Node.init(inType: gReluType) return beginNode diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift index 40069f6550ea00e986926f40c5fc2a2d4bf22a83..126e8bc9545b4bc6b15364096ac18502b0641a70 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift @@ -46,7 +46,7 @@ class ConvAddOp: Operator, ConvAddParam

>, static func fusionNode() -> Node { let beginNode = Node.init(inType: gConvType) _ = beginNode - --> Node.init(inType: gElementwiseAdd) + --> Node.init(inType: gElementwiseAddType) return beginNode } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvBNReluOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvBNReluOp.swift new file mode 100644 index 0000000000000000000000000000000000000000..ef4c0d794e8e0958819d8a423524e08bb281ac84 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvBNReluOp.swift @@ -0,0 +1,129 @@ +/* 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 ConvBNReluParam: OpParam { + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + filter = try ConvBNReluParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) + input = try ConvBNReluParam.input(inputs: opDesc.inputs, from: inScope) + output = try ConvBNReluParam.outputOut(outputs: opDesc.outputs, from: inScope) + stride = try ConvBNReluParam.getAttr(key: "strides", attrs: opDesc.attrs) + paddings = try ConvBNReluParam.getAttr(key: "paddings", attrs: opDesc.attrs) + dilations = try ConvBNReluParam.getAttr(key: "dilations", attrs: opDesc.attrs) + epsilon = try ConvBNReluParam.getAttr(key: "epsilon", attrs: opDesc.attrs) + + groups = try ConvBNReluParam.getAttr(key: "groups", attrs: opDesc.attrs) + variance = try ConvBNReluParam.inputVariance(inputs: opDesc.paraInputs, from: inScope) + bias = try ConvBNReluParam.inputBiase(inputs: opDesc.paraInputs, from: inScope) + scale = try ConvBNReluParam.inputScale(inputs: opDesc.paraInputs, from: inScope) + mean = try ConvBNReluParam.inputMean(inputs: opDesc.paraInputs, from: inScope) + } catch let error { + throw error + } + } + + let input: Texture

+ + let variance: Tensor + let bias: Tensor + let mean: Tensor + let scale: Tensor + let filter: Tensor + let epsilon: Float32 + var newScale: MTLBuffer? + var newBiase: MTLBuffer? + + var output: Texture

+ let stride: [Int32] + let paddings: [Int32] + let dilations: [Int32] + let groups: Int +} + +class ConvBNReluOp: Operator, ConvBNReluParam

>, Runable, Creator, InferShaperable, Fusion{ + typealias OpType = ConvBNReluOp

+ + func inferShape() { + let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations + + var outDim = [inDims[0]] + for i in 0.. Node { + let beginNode = Node.init(inType: gConvType) + _ = beginNode + --> Node.init(inType: gBatchNormType) + --> Node.init(inType: gReluType) + return beginNode + } + + static func change() -> [String : [(from: String, to: String)]] { + return [:] + } + + static func fusionType() -> String { + return gConvBnReluType + } + + func delogOutput() { + + // let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false) + // para.filter.logDataPointer(header: "filter data pointer: ") + // print("filter: \(para.filter)") + + // print("biase: \(para.y)") + // print("padding: \(para.paddings)") + // print("stride: \(para.stride)") + + // let _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false) + // let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false) + // let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false) + + let output = para.output.metalTexture.floatArray { (p: P) -> P in + return p + } + // + writeToLibrary(fileName: "output_112x112x32_2", array: output) + print(" write done") + + // let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false) + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/DepthwiseConvOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/DepthwiseConvOp.swift new file mode 100644 index 0000000000000000000000000000000000000000..272c386cddd52b95f61c2e9aca3aa72cd997cc1e --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/DepthwiseConvOp.swift @@ -0,0 +1,63 @@ +/* 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 DepthConvOp: Operator, ConvParam

>, Runable, Creator, InferShaperable { + required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { + do { + try super.init(device: device, opDesc: opDesc, inScope: inScope) + } catch let error { + throw error + } + } + + func inferShape() { + let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations + + var outDim = [inDims[0]] + for i in 0.. + + func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } + } + + func delogOutput() { + print("conv output : ") + print(para.output.metalTexture) + // let _: Float16? = para.output.metalTexture.logDesc() + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/DwConvBNReluOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/DwConvBNReluOp.swift new file mode 100644 index 0000000000000000000000000000000000000000..da6651d94ff299cbe91aa547b3ae51d68c0aa4ec --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/DwConvBNReluOp.swift @@ -0,0 +1,89 @@ +/* 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 DwConvBNReluOp: Operator, ConvBNReluParam

>, Runable, Creator, InferShaperable, Fusion{ + typealias OpType = ConvBNReluOp

+ + func inferShape() { + let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations + + var outDim = [inDims[0]] + for i in 0.. Node { + let beginNode = Node.init(inType: gDepthConvType) + _ = beginNode + --> Node.init(inType: gBatchNormType) + --> Node.init(inType: gReluType) + return beginNode + } + + static func change() -> [String : [(from: String, to: String)]] { + return [:] + } + + static func fusionType() -> String { + return gDwConvBnReluType + } + + func delogOutput() { + + // let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false) + // para.filter.logDataPointer(header: "filter data pointer: ") + // print("filter: \(para.filter)") + + // print("biase: \(para.y)") + // print("padding: \(para.paddings)") + // print("stride: \(para.stride)") + + // let _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false) + // let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false) + // let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false) + + let output = para.output.metalTexture.floatArray { (p: P) -> P in + return p + } + // + writeToLibrary(fileName: "output_112x112x32_2", array: output) + print(" write done") + + // let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false) + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/BoxcoderKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/BoxcoderKernel.swift new file mode 100644 index 0000000000000000000000000000000000000000..6cc15628cb8bddf28730d88408358d37a94bed93 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/BoxcoderKernel.swift @@ -0,0 +1,31 @@ +/* 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 BoxcoderKernel: Kernel, Computable{ + func compute(commandBuffer: MTLCommandBuffer, param: BoxcoderParam

) 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) + encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() + } + + required init(device: MTLDevice, param: BoxcoderParam

) { + super.init(device: device, inFunctionName: "priorbox") + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConcatKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConcatKernel.swift new file mode 100644 index 0000000000000000000000000000000000000000..c1cc11d16f28a5d1b91fc2ce9670b04492f25ec8 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConcatKernel.swift @@ -0,0 +1,31 @@ +/* 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 ConcatKernel: Kernel, Computable{ + func compute(commandBuffer: MTLCommandBuffer, param: ConcatParam

) 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) + encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() + } + + required init(device: MTLDevice, param: ConcatParam

) { + super.init(device: device, inFunctionName: "priorbox") + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift new file mode 100644 index 0000000000000000000000000000000000000000..867787b10cf5527e4f52d2f904b8cbc1b7f626f8 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift @@ -0,0 +1,137 @@ +/* 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 ConvBNReluTestParam: TestParam { + let inputTexture: MTLTexture + let outputTexture: MTLTexture + var metalParam: MetalConvParam + let filterBuffer: MTLBuffer + let biaseBuffer: MTLBuffer + let newScaleBuffer: MTLBuffer + let newBiaseBuffer: MTLBuffer + let filterSize: (width: Int, height: Int, channel: Int) + init(inInputTexture: MTLTexture, inOutputTexture: MTLTexture, inMetalParam: MetalConvParam, inFilterBuffer: MTLBuffer, inBiaseBuffer: MTLBuffer, inNewScaleBuffer: MTLBuffer, inNewBiaseBuffer: MTLBuffer, inFilterSize: (width: Int, height: Int, channel: Int)) { + inputTexture = inInputTexture + outputTexture = inOutputTexture + metalParam = inMetalParam + filterBuffer = inFilterBuffer + biaseBuffer = inBiaseBuffer + newScaleBuffer = inNewScaleBuffer + newBiaseBuffer = inNewBiaseBuffer + filterSize = inFilterSize + } +} + +class ConvBNReluKernel: Kernel, Computable, Testable { + required init(device: MTLDevice, testParam: ConvBNReluTestParam) { + if testParam.filterSize.width == 1 && testParam.filterSize.height == 1 { + super.init(device: device, inFunctionName: "conv_add_batch_norm_relu_1x1") + } else if testParam.filterSize.channel == 1 { + super.init(device: device, inFunctionName: "depthwise_conv_add_batch_norm_relu_3x3") + } else { + super.init(device: device, inFunctionName: "conv_add_batch_norm_relu_3x3") + } + } + + var metalParam: MetalConvParam! + + required init(device: MTLDevice, param: ConvBNReluParam

) { + + if param.filter.width == 1 && param.filter.height == 1 { + super.init(device: device, inFunctionName: "conv_add_batch_norm_relu_1x1") + } else if param.filter.channel == 1 { + super.init(device: device, inFunctionName: "depthwise_conv_add_batch_norm_relu_3x3") + } else { + super.init(device: device, inFunctionName: "conv_add_batch_norm_relu_3x3") + } + + param.filter.initBuffer(device: device, precision: Tensor.BufferPrecision.Float32) + + param.variance.initBuffer(device: device) + param.mean.initBuffer(device: device) + param.scale.initBuffer(device: device) + param.bias.initBuffer(device: device) + + let offsetX = param.filter.width/2 - Int(param.paddings[0]) + let offsetY = param.filter.height/2 - Int(param.paddings[1]) + + print("offset x: \(offsetX)") + print("offset y: \(offsetY)") + + let offsetZ = 0.0 + metalParam = MetalConvParam.init(offsetX: Int16(offsetX), offsetY: Int16(offsetY), offsetZ: Int16(offsetZ), strideX: UInt16(param.stride[0]), strideY: UInt16(param.stride[1]), paddedZ: UInt16(param.input.metalTexture.arrayLength * 4 - param.input.dim[3])) + + var invs: [P] = [] + let varianceContents = param.variance.buffer.contents().assumingMemoryBound(to: P.self) + + for i in 0...stride { + let inv = 1.0/pow(Float32.init(varianceContents[i]) + param.epsilon, 0.5) + invs.append(P(inv)) + } + + let newScale: UnsafeMutablePointer

= UnsafeMutablePointer

.allocate(capacity: param.scale.buffer.length) + let newBiase: UnsafeMutablePointer

= UnsafeMutablePointer

.allocate(capacity: param.bias.buffer.length) + + let scaleContents = param.scale.buffer.contents().assumingMemoryBound(to: P.self) + let biaseContents = param.bias.buffer.contents().assumingMemoryBound(to: P.self) + let meanContents = param.mean.buffer.contents().assumingMemoryBound(to: P.self) + for i in 0...stride { + newScale[i] = invs[i] * scaleContents[i] + newBiase[i] = biaseContents[i] - meanContents[i] * invs[i] * scaleContents[i] + } + + param.newBiase = device.makeBuffer(bytes: newBiase, length: param.bias.buffer.length) + param.newScale = device.makeBuffer(bytes: newScale, length: param.scale.buffer.length) + + newScale.deinitialize(count: param.scale.buffer.length) + newScale.deallocate() + + newBiase.deinitialize(count: param.bias.buffer.length) + newBiase.deallocate() + } + + func compute(commandBuffer: MTLCommandBuffer, param: ConvBNReluParam

) 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) + encoder.setBytes(&metalParam, length: MemoryLayout.size, index: 0) + encoder.setBuffer(param.filter.buffer, offset: 0, index: 1) + encoder.setBuffer(param.newScale!, offset: 0, index: 3) + encoder.setBuffer(param.newBiase!, offset: 0, index: 4) + encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() + } + + public func test(commandBuffer: MTLCommandBuffer, param: ConvBNReluTestParam) { + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { + fatalError() + } + + encoder.setTexture(param.inputTexture, index: 0) + encoder.setTexture(param.outputTexture, index: 1) + var inMetalParam = param.metalParam + encoder.setBytes(&inMetalParam, length: MemoryLayout.size, index: 0) + encoder.setBuffer(param.filterBuffer, offset: 0, index: 1) + encoder.setBuffer(param.biaseBuffer, offset: 0, index: 2) + encoder.setBuffer(param.newScaleBuffer, offset: 0, index: 3) + encoder.setBuffer(param.newBiaseBuffer, offset: 0, index: 4) + encoder.dispatch(computePipline: pipline, outTexture: param.outputTexture) + encoder.endEncoding() + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/MulticlassNMSKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/MulticlassNMSKernel.swift new file mode 100644 index 0000000000000000000000000000000000000000..a85931af5e6ef042af62f754181de2beccb740bb --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/MulticlassNMSKernel.swift @@ -0,0 +1,31 @@ +/* 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 MulticlassNMSKernel: Kernel, Computable{ + func compute(commandBuffer: MTLCommandBuffer, param: MulticlassNMSParam

) 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) + encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() + } + + required init(device: MTLDevice, param: MulticlassNMSParam

) { + super.init(device: device, inFunctionName: "priorbox") + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift new file mode 100644 index 0000000000000000000000000000000000000000..8e33c19f77675b446b7b08d70475831cfea7d989 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift @@ -0,0 +1,31 @@ +/* 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 PriorBoxKernel: Kernel, Computable{ + func compute(commandBuffer: MTLCommandBuffer, param: PriorBoxParam

) 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) + encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() + } + + required init(device: MTLDevice, param: PriorBoxParam

) { + super.init(device: device, inFunctionName: "priorbox") + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift new file mode 100644 index 0000000000000000000000000000000000000000..c3868dd0708a9859f3133c7b71461a64f194c079 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift @@ -0,0 +1,31 @@ +/* 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 TransposeKernel: Kernel, Computable{ + func compute(commandBuffer: MTLCommandBuffer, param: TransposeParam

) 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) + encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() + } + + required init(device: MTLDevice, param: TransposeParam

) { + super.init(device: device, inFunctionName: "priorbox") + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift new file mode 100644 index 0000000000000000000000000000000000000000..dae584d52c6d2335908d28741c4171c8699eff80 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift @@ -0,0 +1,47 @@ +///* 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 MulticlassNMSParam: OpParam { + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + fatalError() + } catch let error { + throw error + } + } + let input: Texture

+ var output: Texture

+} + +class MulticlassNMSOp: Operator, MulticlassNMSParam

>, Runable, Creator, InferShaperable{ + + func inferShape() { + para.output.dim = para.input.dim + } + + typealias OpType = MulticlassNMSOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } + } +} + + + diff --git a/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift new file mode 100644 index 0000000000000000000000000000000000000000..772ed0e2f8a558eb8d54c1eacc0dd695c82511f2 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift @@ -0,0 +1,47 @@ +///* 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 PriorBoxParam: OpParam { + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + fatalError() + } catch let error { + throw error + } + } + let input: Texture

+ var output: Texture

+} + +class PriorBoxOp: Operator, PriorBoxParam

>, Runable, Creator, InferShaperable{ + + func inferShape() { + para.output.dim = para.input.dim + } + + typealias OpType = PriorBoxOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } + } +} + + + diff --git a/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift new file mode 100644 index 0000000000000000000000000000000000000000..f9a249484acb8533b66ed56635bc8e728c9ba3dd --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift @@ -0,0 +1,47 @@ +///* 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 TransposeParam: OpParam { + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + fatalError() + } catch let error { + throw error + } + } + let input: Texture

+ var output: Texture

+} + +class TransposeOp: Operator, TransposeParam

>, Runable, Creator, InferShaperable{ + + func inferShape() { + para.output.dim = para.input.dim + } + + typealias OpType = TransposeOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } + } +} + + + diff --git a/metal/paddle-mobile/paddle-mobile/Program/ProgramOptimize.swift b/metal/paddle-mobile/paddle-mobile/Program/ProgramOptimize.swift index d819cdad533e444c327e95baff7bf87e902d6bff..e744901a5ce7a0ac2363336fababc751024abd61 100644 --- a/metal/paddle-mobile/paddle-mobile/Program/ProgramOptimize.swift +++ b/metal/paddle-mobile/paddle-mobile/Program/ProgramOptimize.swift @@ -143,7 +143,12 @@ extension Node: Equatable { } class ProgramOptimize { - let fusionOps: [Fusion.Type] = [ConvAddBatchNormReluOp

.self, ConvAddOp

.self] + // register fusion + let fusionOps: [Fusion.Type] = [ConvAddBatchNormReluOp

.self, + ConvAddOp

.self, + ConvBNReluOp

.self, + DwConvBNReluOp

.self] + func optimize(originProgramDesc: ProgramDesc) -> ProgramDesc { guard originProgramDesc.blocks.count == 1 else { diff --git a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift index 81894664c5dc4acb1a5edd4485543bb20a285ea4..d217d74abfd3b763d7b923066608ad098a5a6884 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift @@ -22,7 +22,6 @@ class InputTexture { mtlTexture = inMTLTexture expectDim = inExpectDim } - } extension InputTexture { @@ -54,19 +53,34 @@ public class Texture: Tensorial { } else if inDim.cout() == 4 { tmpTextureDes.height = inDim[1] tmpTextureDes.width = inDim[2] -// print("n : \(inDim[0])") -// print(inDim[3] * inDim[0]) tmpTextureDes.depth = 1 tmpTextureDes.arrayLength = (inDim[3] * inDim[0] + 3)/4 tmpTextureDes.textureType = .type2DArray } else if inDim.cout() == 2 { - tmpTextureDes.height = 1 - tmpTextureDes.width = 1 +// tmpTextureDes.height = 1 +// tmpTextureDes.width = 1 +// tmpTextureDes.depth = 1 +// tmpTextureDes.arrayLength = (inDim[0] * inDim[1] + 3)/4 + + tmpTextureDes.width = inDim[0] + tmpTextureDes.height = inDim[1] tmpTextureDes.depth = 1 - tmpTextureDes.arrayLength = (inDim[0] * inDim[1] + 3)/4 + tmpTextureDes.arrayLength = 1 tmpTextureDes.textureType = .type2DArray } else { - fatalError(" not suuprt ") + /* + var name: box_coder_0.tmp_0 + in var tensor desc dims size: 3 + var tensor desc dim 0 value: -1 + var tensor desc dim 1 value: 1917 + var tensor desc dim 2 value: 4 + */ + + tmpTextureDes.height = inDim[1] + tmpTextureDes.width = inDim[2] + tmpTextureDes.depth = 1 + tmpTextureDes.arrayLength = 1 + tmpTextureDes.textureType = .type2DArray } if MemoryLayout

.size == 1 { @@ -79,7 +93,7 @@ public class Texture: Tensorial { } // tmpTextureDes.pixelFormat = .rgba16Float - + tmpTextureDes.usage = [.shaderRead, .shaderWrite] tmpTextureDes.storageMode = .shared textureDesc = tmpTextureDes