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 9f80a84acc4f475730cf686bd018922ecb769f4e..4036ea7722f78ca1e14ac79dc496b0f437948df1 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj +++ b/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj @@ -14,10 +14,10 @@ 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 */; }; - FC27991021341CE5000B6BAD /* Net.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC27990F21341CE5000B6BAD /* Net.swift */; }; FC27991321343A3A000B6BAD /* CPUCompute.mm in Sources */ = {isa = PBXBuildFile; fileRef = FC27991221343A3A000B6BAD /* CPUCompute.mm */; }; FC3C800F2133F46600D1295E /* MobileNetSSD.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3C800E2133F46600D1295E /* MobileNetSSD.swift */; }; FC3C80112133F4AB00D1295E /* MobileNet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3C80102133F4AB00D1295E /* MobileNet.swift */; }; + FC4FD95121402B610073E130 /* PaddleMobile.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC4FD95021402B610073E130 /* PaddleMobile.swift */; }; FC8CFEE2213524EA0094D569 /* Genet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC8CFEE1213524EA0094D569 /* Genet.swift */; }; FC8CFEE62135452C0094D569 /* genet_params in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEE42135452B0094D569 /* genet_params */; }; FC8CFEE72135452C0094D569 /* genet_model in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEE52135452B0094D569 /* genet_model */; }; @@ -61,12 +61,12 @@ 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 = ""; }; - FC27990F21341CE5000B6BAD /* Net.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Net.swift; sourceTree = ""; }; FC27991121343A39000B6BAD /* paddle-mobile-demo-Bridging-Header.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = "paddle-mobile-demo-Bridging-Header.h"; sourceTree = ""; }; FC27991221343A3A000B6BAD /* CPUCompute.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = CPUCompute.mm; sourceTree = ""; }; FC27991421343A46000B6BAD /* CPUCompute.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = CPUCompute.h; sourceTree = ""; }; FC3C800E2133F46600D1295E /* MobileNetSSD.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNetSSD.swift; sourceTree = ""; }; FC3C80102133F4AB00D1295E /* MobileNet.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNet.swift; sourceTree = ""; }; + FC4FD95021402B610073E130 /* PaddleMobile.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; name = PaddleMobile.swift; path = ../../../../../../Desktop/PaddleMobile.swift; sourceTree = ""; }; FC8CFEE1213524EA0094D569 /* Genet.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Genet.swift; sourceTree = ""; }; FC8CFEE42135452B0094D569 /* genet_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = genet_params; sourceTree = ""; }; FC8CFEE52135452B0094D569 /* genet_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = genet_model; sourceTree = ""; }; @@ -176,11 +176,11 @@ FC8CFED2213519540094D569 /* Net */ = { isa = PBXGroup; children = ( + FC4FD95021402B610073E130 /* PaddleMobile.swift */, FC013927210204A3008100E3 /* PreProcessKernel.metal */, FCBCCC542122EF5400D94F7E /* MetalHelper.swift */, FC3C800E2133F46600D1295E /* MobileNetSSD.swift */, FC3C80102133F4AB00D1295E /* MobileNet.swift */, - FC27990F21341CE5000B6BAD /* Net.swift */, FC27991221343A3A000B6BAD /* CPUCompute.mm */, FC27991421343A46000B6BAD /* CPUCompute.h */, FC8CFEE1213524EA0094D569 /* Genet.swift */, @@ -342,11 +342,11 @@ files = ( FC039B8420E11C550081E9F8 /* ViewController.swift in Sources */, FC013928210204A3008100E3 /* PreProcessKernel.metal in Sources */, - FC27991021341CE5000B6BAD /* Net.swift in Sources */, FC8CFEE2213524EA0094D569 /* Genet.swift in Sources */, FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */, FC27991321343A3A000B6BAD /* CPUCompute.mm in Sources */, FC3C80112133F4AB00D1295E /* MobileNet.swift in Sources */, + FC4FD95121402B610073E130 /* PaddleMobile.swift in Sources */, FC3C800F2133F46600D1295E /* MobileNetSSD.swift in Sources */, FC039B8220E11C550081E9F8 /* AppDelegate.swift in Sources */, ); @@ -494,19 +494,19 @@ buildSettings = { ASSETCATALOG_COMPILER_APPICON_NAME = AppIcon; CLANG_ENABLE_MODULES = YES; - CODE_SIGN_IDENTITY = "iPhone Distribution"; - CODE_SIGN_STYLE = Manual; - DEVELOPMENT_TEAM = 6T9LLJKSM4; + CODE_SIGN_IDENTITY = "iPhone Developer"; + CODE_SIGN_STYLE = Automatic; + DEVELOPMENT_TEAM = A798K58VVL; INFOPLIST_FILE = "paddle-mobile-demo/Info.plist"; IPHONEOS_DEPLOYMENT_TARGET = 9.0; LD_RUNPATH_SEARCH_PATHS = ( "$(inherited)", "@executable_path/Frameworks", ); - PRODUCT_BUNDLE_IDENTIFIER = com.baidu.mms.qa; + PRODUCT_BUNDLE_IDENTIFIER = "com.baidu.paddle-mobile"; PRODUCT_NAME = "$(TARGET_NAME)"; - PROVISIONING_PROFILE = "ba9c4b24-7bd0-49c5-93cd-e3051e775d6c"; - PROVISIONING_PROFILE_SPECIFIER = Distribution_MMS; + PROVISIONING_PROFILE = ""; + PROVISIONING_PROFILE_SPECIFIER = ""; SWIFT_OBJC_BRIDGING_HEADER = "paddle-mobile-demo/paddle-mobile-demo-Bridging-Header.h"; SWIFT_OPTIMIZATION_LEVEL = "-Onone"; SWIFT_VERSION = 4.0; @@ -520,19 +520,19 @@ buildSettings = { ASSETCATALOG_COMPILER_APPICON_NAME = AppIcon; CLANG_ENABLE_MODULES = YES; - CODE_SIGN_IDENTITY = "iPhone Distribution"; - CODE_SIGN_STYLE = Manual; - DEVELOPMENT_TEAM = 6T9LLJKSM4; + CODE_SIGN_IDENTITY = "iPhone Developer"; + CODE_SIGN_STYLE = Automatic; + DEVELOPMENT_TEAM = A798K58VVL; INFOPLIST_FILE = "paddle-mobile-demo/Info.plist"; IPHONEOS_DEPLOYMENT_TARGET = 9.0; LD_RUNPATH_SEARCH_PATHS = ( "$(inherited)", "@executable_path/Frameworks", ); - PRODUCT_BUNDLE_IDENTIFIER = com.baidu.mms.qa; + PRODUCT_BUNDLE_IDENTIFIER = "com.baidu.paddle-mobile"; PRODUCT_NAME = "$(TARGET_NAME)"; - PROVISIONING_PROFILE = "ba9c4b24-7bd0-49c5-93cd-e3051e775d6c"; - PROVISIONING_PROFILE_SPECIFIER = Distribution_MMS; + PROVISIONING_PROFILE = ""; + PROVISIONING_PROFILE_SPECIFIER = ""; SWIFT_OBJC_BRIDGING_HEADER = "paddle-mobile-demo/paddle-mobile-demo-Bridging-Header.h"; SWIFT_VERSION = 4.0; TARGETED_DEVICE_FAMILY = "1,2"; diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/Genet.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/Genet.swift index 8b1d2c4ce4eb51e3fc8ca578b72fd3fd776faa89..9f1a31920de671350649091e184dc5d452f331a6 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/Genet.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/Genet.swift @@ -17,10 +17,6 @@ import paddle_mobile class Genet: Net { - var program: Program? - - var executor: Executor? - let except: Int = 0 class GenetPreProccess: CusomKernel { diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MetalHelper.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MetalHelper.swift index e7c2cccea832d33847669d8c1aeeeed4a47ab0f8..d314e8b3f8845ef95b36b4b25e61809d353f0f24 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MetalHelper.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MetalHelper.swift @@ -28,25 +28,6 @@ class MetalHelper { textureLoader = MTKTextureLoader.init(device: device) } - static func scaleTexture(queue: MTLCommandQueue, input: MTLTexture, size:(width: Int, height: Int), complete: @escaping (MTLTexture) -> Void) { - guard let buffer = queue.makeCommandBuffer() else { - fatalError() - } - - let scaleKernel = ScaleKernel.init(device: MetalHelper.shared.device, shape: CusomKernel.Shape.init(inWidth: size.width, inHeight: size.height, inChannel: 3)) - - do { - try scaleKernel.compute(inputTexuture: input, commandBuffer: buffer) - } catch let error { - print(error) - fatalError() - } - - buffer.addCompletedHandler { (buffer) in - complete(scaleKernel.outputTexture) - } - buffer.commit() - } } diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNet.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNet.swift index 3c06c108b9be77eb6f58e17dcc95ce3c96891b09..4579ee07356c3c599f58060c9dc11d8401992817 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNet.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNet.swift @@ -15,12 +15,10 @@ import Foundation import paddle_mobile + + class MobileNet: Net{ - var program: Program? - - var executor: Executor? - let except: Int = 0 class MobilenetPreProccess: CusomKernel { diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNetSSD.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNetSSD.swift index 1349a93a4724e1ada91ea4d63638d56a55fcdb25..83767f0bc3ebb3def795dd45643a4a77d8f59dd1 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNetSSD.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNetSSD.swift @@ -14,13 +14,8 @@ import Foundation import paddle_mobile -//import class MobileNet_ssd_hand: Net{ - - var program: Program? - var executor: Executor? - let except: Int = 2 class MobilenetssdPreProccess: CusomKernel { init(device: MTLDevice) { @@ -83,36 +78,11 @@ class MobileNet_ssd_hand: Net{ let modelDir: String - -// let paramPointer: UnsafeMutableRawPointer -// -// let paramSize: Int -// -// let modelPointer: UnsafeMutableRawPointer -// -// let modelSize: Int -// -// /** -// * inParamPointer: 参数文件内存地址 -// * inParamSize: 参数文件大小(字节数) -// * inModelPointer: 模型文件内存地址 -// * inModelSize: 模型文件大小(字节数) -// */ -// init(inParamPointer: UnsafeMutableRawPointer, inParamSize: Int, inModelPointer: UnsafeMutableRawPointer, inModelSize: Int) { -// paramPointer = inParamPointer -// paramSize = inParamSize -// modelPointer = inModelPointer -// modelSize = inModelSize -//// fatalError() -// } - - 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) -// fatalError() } } diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/Net.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/Net.swift deleted file mode 100644 index c643b4b63d263b2020170483c955bea4d6b6e404..0000000000000000000000000000000000000000 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/Net.swift +++ /dev/null @@ -1,89 +0,0 @@ -/* 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 - -import UIKit -import MetalKit -import Foundation -import paddle_mobile -import MetalPerformanceShaders - -class ScaleKernel: CusomKernel { - init(device: MTLDevice, shape: Shape) { - super.init(device: device, inFunctionName: "scale", outputDim: shape, usePaddleMobileLib: false) - } -} - -protocol Net { - var program: Program? { get set } - var executor: Executor? { get set } - var except: Int { get } - var dim: (n: Int, h: Int, w: Int, c: Int) { get } - var modelPath: String { get } - var paramPath: String { get } - var modelDir: String { get } - var preprocessKernel: CusomKernel { get } - func getTexture(image: CGImage, getTexture: @escaping (MTLTexture) -> Void) - func resultStr(res: [Float]) -> String - func fetchResult(paddleMobileRes: ResultHolder) -> [Float32] - mutating func load() throws - - func predict(inTexture: MTLTexture, completion: @escaping ((time:TimeInterval, resultArray: [Float32])) -> Void) throws - mutating func clear() -} - -extension Net { - mutating func load() throws { - let queue = MetalHelper.shared.queue - let loader = Loader.init() - do { - program = try loader.load(device: MetalHelper.shared.device, modelPath: modelPath, paraPath: paramPath) - executor = try Executor.init(inDevice: MetalHelper.shared.device, inQueue: queue, inProgram: program!) - } catch let error { - throw error - } - } - - func predict(inTexture: MTLTexture, completion: @escaping ((time:TimeInterval, resultArray: [Float32])) -> Void) throws { - guard let inExecutor = executor else { - fatalError(" 请先 load ") - } - try inExecutor.predict(input: inTexture, dim: [dim.n, dim.h, dim.w, dim.c], completionHandle: { (result) in - - var resultArr:[Float32] = [] - resultArr = self.fetchResult(paddleMobileRes: result) - completion((time: TimeInterval(result.elapsedTime), resultArray: resultArr)) - - }, preProcessKernle: preprocessKernel, except: except) - } - - mutating func clear() { - executor?.clear() - program = nil - executor = nil - } - - 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: (dim.w, dim.h)) { (resTexture) in - getTexture(resTexture) - } - } - - func fetchResult(paddleMobileRes: ResultHolder) -> [Float32] { - return paddleMobileRes.resultArr - } - -} diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/PaddleMobile.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/PaddleMobile.swift new file mode 100644 index 0000000000000000000000000000000000000000..a954328acae3a80643ad849d58cd6ac86bf7865e --- /dev/null +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/PaddleMobile.swift @@ -0,0 +1,9 @@ +// +// PaddleMobile.swift +// paddle-mobile-demo +// +// Created by liuRuiLong on 2018/9/5. +// Copyright © 2018年 orange. All rights reserved. +// + +import Foundation diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift index 47708a70c9f3c69a42c5ab220180ad6a4f695556..cecf14411634e93c44ee7624ee1872dc2539938d 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift @@ -25,7 +25,7 @@ let modelHelperMap: [SupportModel : Net] = [.mobilenet_ssd : MobileNet_ssd_hand. //let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()] enum SupportModel: String{ -// case mobilenet = "mobilenet" + // case mobilenet = "mobilenet" case mobilenet_ssd = "mobilenetssd" case genet = "genet" static func supportedModels() -> [SupportModel] { @@ -40,6 +40,7 @@ class ViewController: UIViewController { @IBOutlet weak var elapsedTimeLabel: UILabel! @IBOutlet weak var modelPickerView: UIPickerView! @IBOutlet weak var threadPickerView: UIPickerView! + var runnner: Runner! var selectImage: UIImage? var modelType: SupportModel = SupportModel.supportedModels()[0] var toPredictTexture: MTLTexture? @@ -56,10 +57,10 @@ class ViewController: UIViewController { @IBAction func loadAct(_ sender: Any) { - do { - try self.net.load() - } catch let error { - print(error) + if runnner.load() { + print(" load success ! ") + } else { + print(" load error ! ") } } @@ -71,7 +72,7 @@ class ViewController: UIViewController { } @IBAction func clearAct(_ sender: Any) { - net.clear() + runnner.clear() } @IBAction func predictAct(_ sender: Any) { @@ -79,26 +80,22 @@ class ViewController: UIViewController { resultTextView.text = "请选择图片 ! " return } - do { - let max = 50 - let startDate = Date.init() - for i in 0.. + let filterSize: (width: Int, height: Int, channel: Int) + init(inMetalParam: MetalConvParam, inFilter: [Float], inBiase: [Float], inFilterSize: (width: Int, height: Int, channel: Int)) { + metalParam = inMetalParam + filterPointer = UnsafeMutableRawPointer.init(mutating: inFilter) + biasePointer = UnsafeMutablePointer.init(mutating: inBiase) + filterSize = inFilterSize + } +} + +@available(iOS 10.0, *) +class CNNMPSConvOp: Operator, CNNConvParam

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

+ + required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { + fatalError() + } + + func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } + } + + func delogOutput() { + } + + static func fusionNode() -> Node { + let beginNode = Node.init(inType: gConvType) + _ = beginNode-->Node.init(inType: gElementwiseAdd); + return beginNode + } + + static func change() -> [String : [(from: String, to: String)]] { + return [:] + } + + static func fusionType() -> String { + return gMPSCNNConvType + } + 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..: BatchNormParam

{ + +} + +class BatchNormReluKernel: Kernel, Computable{ + + + typealias ParamType = BatchNormReluParam

+ var newScale: MTLBuffer + var newBias: MTLBuffer + + required init(device: MTLDevice, testParam: BatchNormReluTestParam) { + + newScale = testParam.newScaleBuffer + newBias = testParam.newBiaseBuffer + + super.init(device: device, inFunctionName: "batch_norm_relu_3x3") + } + + required init(device: MTLDevice, param: BatchNormReluParam

) { + guard let newScale = device.makeBuffer(length: param.inputScale.buffer.length) else { + fatalError() + } + guard let newBias = device.makeBuffer(length: param.inputBias.buffer.length) else { + fatalError() + } + self.newScale = newScale + self.newBias = newBias + + super.init(device: device, inFunctionName: "batch_norm_relu_3x3") + + + let varianceBuffer : MTLBuffer = param.inputVariance.buffer + + var invStd: [Float32] = Array(repeating: 0, count: varianceBuffer.length) + let varianceContents = varianceBuffer.contents().assumingMemoryBound(to: P.self) + for i in 0..<(varianceBuffer.length / MemoryLayout

.stride) { + invStd[i] = 1 / (Float32(varianceContents[i]) + param.epsilon).squareRoot() + } + + let newScaleContents = newScale.contents().assumingMemoryBound(to: P.self) + let newBiasContents = newBias.contents().assumingMemoryBound(to: P.self) + let scale : MTLBuffer = param.inputScale.buffer + let scaleContents = scale.contents().assumingMemoryBound(to: P.self) + let bias : MTLBuffer = param.inputBias.buffer + let biasContents = bias.contents().assumingMemoryBound(to: P.self) + let meanContents = param.inputMean.buffer.contents().assumingMemoryBound(to: P.self) + + for i in 0..<(newScale.length / MemoryLayout

.stride) { + newScaleContents[i] = P(invStd[i] * Float32(scaleContents[i])) + newBiasContents[i] = P(Float32(biasContents[i]) - Float32(meanContents[i]) * invStd[i] * Float32(scaleContents[i])) + } + } + + func compute(commandBuffer: MTLCommandBuffer, param: BatchNormReluParam

) throws { + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { + fatalError() + } + encoder.setTexture(param.input as? MTLTexture, index: 0) + encoder.setTexture(param.output as? MTLTexture, index: 1) + encoder.setBuffer(newScale, offset: 0, index: 1) + encoder.setBuffer(newBias, offset: 0, index: 1) + encoder.dispatch(computePipline: pipline, outTexture: param.output as! MTLTexture) + encoder.endEncoding() + } + + func testCompute(commandBuffer: MTLCommandBuffer, testParam: BatchNormReluTestParam) throws { + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { + fatalError() + } + encoder.setTexture(testParam.inputTexture, index: 0) + encoder.setTexture(testParam.outputTexture, index: 1) + encoder.setBuffer(newScale, offset: 0, index: 0) + encoder.setBuffer(newBias, offset: 0, index: 1) + encoder.dispatch(computePipline: pipline, outTexture: testParam.outputTexture) + encoder.endEncoding() + } + + +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/CNNConvKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/CNNConvKernel.swift new file mode 100644 index 0000000000000000000000000000000000000000..14a5bd521455632c8a67e4c1a8ebdedc6c460aa5 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/CNNConvKernel.swift @@ -0,0 +1,176 @@ +// +// CNNConvKernel.swift +// paddle-mobile +// + +import Foundation +import Metal +import Accelerate +import MetalPerformanceShaders + +@available(iOS 10.0, *) +class WeightsDataSource: NSObject, MPSCNNConvolutionDataSource { + + let desc: MPSCNNConvolutionDescriptor + let weight:UnsafeMutableRawPointer + let bias:UnsafeMutablePointer + + + + init(inDesc: MPSCNNConvolutionDescriptor, inWeight: UnsafeMutableRawPointer, inBias: UnsafeMutablePointer) { + desc = inDesc + weight = inWeight + bias = inBias + } + + + func dataType() -> MPSDataType { + return .float32 + } + + func descriptor() -> MPSCNNConvolutionDescriptor { + return desc + } + + func weights() -> UnsafeMutableRawPointer { + return self.weight + } + + func biasTerms() -> UnsafeMutablePointer? { + return self.bias + } + + func load() -> Bool { + return true + } + + func purge() { + } + + func label() -> String? { + return "Conv" + } + + +} + +@available(iOS 10.0, *) +class CNNConvParam: OpParam{ + + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + filter = try CNNConvParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) + input = try CNNConvParam.input(inputs: opDesc.inputs, from: inScope) + output = try CNNConvParam.outputOut(outputs: opDesc.outputs, from: inScope) + stride = try CNNConvParam.getAttr(key: "strides", attrs: opDesc.attrs) + paddings = try CNNConvParam.getAttr(key: "paddings", attrs: opDesc.attrs) + // 暂时不用关心 + dilations = try CNNConvParam.getAttr(key: "dilations", attrs: opDesc.attrs) + // 暂时不用关心 + groups = try CNNConvParam.getAttr(key: "groups", attrs: opDesc.attrs) + + variance = try CNNConvParam.inputVariance(inputs: opDesc.paraInputs, from: inScope) + // bias + y = try CNNConvParam.inputY(inputs: opDesc.paraInputs, from: inScope) + } catch let error { + throw error + } + } + + var input: Texture

+ let variance: Tensor + let y: Tensor + let filter: Tensor + var output: Texture

+ let stride: [Int32] + let paddings: [Int32] + let dilations: [Int32] + let groups: Int +} + +@available(iOS 10.0, *) +class CNNConvKernel: Kernel, Computable { + + typealias ParamType = CNNConvParam

+ + var mpsImageCreator: MpsImageCreator

? + var activation:MPSCNNNeuron? + var conv:MPSCNNConvolution? + var weightDataSource:WeightsDataSource? + var param: CNNConvParam

? + var device: MTLDevice? + + + required init(device:MTLDevice, testParam:CNNMPSConvTestParam) { + self.device = device + + let desc = MPSCNNConvolutionDescriptor(kernelWidth: testParam.filterSize.width, kernelHeight: testParam.filterSize.height, inputFeatureChannels: testParam.filterSize.channel, outputFeatureChannels: testParam.filterSize.channel, neuronFilter: activation) + + desc.strideInPixelsX = Int(testParam.metalParam.offsetX) + desc.strideInPixelsY = Int(testParam.metalParam.offsetY) + + + weightDataSource = WeightsDataSource(inDesc: desc, inWeight:testParam.filterPointer, inBias:testParam.biasePointer) + + if #available(iOS 11.0, *) { + conv = MPSCNNConvolution(device: self.device!, weights: weightDataSource!) + } else { + // Fallback on earlier versions + } + + super.init(device: device, inFunctionName: "") + } + + required init(device:MTLDevice, param:CNNConvParam

) { + + self.device = device + + let inChannels: Int + let outChannels: Int + + if param.y.dim.cout() == 4 { + inChannels = (param.y.dim[3]) + outChannels = inChannels + } else { + inChannels = 0 + outChannels = inChannels + } + + let desc = MPSCNNConvolutionDescriptor(kernelWidth: param.filter.width, kernelHeight: param.filter.height, inputFeatureChannels: inChannels, outputFeatureChannels: outChannels, neuronFilter: activation) + + desc.strideInPixelsX = Int(param.stride[0]) + desc.strideInPixelsY = Int(param.stride[1]) + + + weightDataSource = WeightsDataSource(inDesc: desc, inWeight:param.filter.data.pointer as! UnsafeMutablePointer, inBias: param.y.data.pointer as! UnsafeMutablePointer) + + if #available(iOS 11.0, *) { + conv = MPSCNNConvolution(device: self.device!, weights: weightDataSource!) + } else { + // Fallback on earlier versions + } + + super.init(device: device, inFunctionName: "") + } + + func compute(commandBuffer: MTLCommandBuffer, param: CNNConvParam

) throws { + let inputImage:MPSImage = (mpsImageCreator?.createMPSImage(device: device!))! + var outputImage = (mpsImageCreator?.createMPSImage(device: device!))! + + // 运算conv和add两个步骤，add用了bias偏差做为参数，被Metal API进行调用 + conv?.encode(commandBuffer: commandBuffer, sourceImage: inputImage, destinationImage: outputImage) + + param.input = outputImage.texture as! Texture

+ } + + func testCompute(commandBuffer: MTLCommandBuffer, testParam: CNNMPSConvTestParam) throws { + let inputImage:MPSImage = (mpsImageCreator?.createMPSImage(device: device!))! + var outputImage = (mpsImageCreator?.createMPSImage(device: device!))! + + // 运算conv和add两个步骤，add用了bias偏差做为参数，被Metal API进行调用 + conv?.encode(commandBuffer: commandBuffer, sourceImage: inputImage, destinationImage: outputImage) + + testParam.outputTexture = outputImage.texture + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/BatchNormRelu.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/BatchNormRelu.metal new file mode 100644 index 0000000000000000000000000000000000000000..eb94408c8ac664be5cf62bc28bfb02825856ebd4 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/BatchNormRelu.metal @@ -0,0 +1,36 @@ +// +// BatchNormRelu.metal +// paddle-mobile +// + +#include +using namespace metal; + +struct MetalConvParam { + short offsetX; + short offsetY; + short offsetZ; + ushort strideX; + ushort strideY; +}; + +kernel void batch_norm_relu_3x3(texture2d_array inTexture [[texture(0)]], + texture2d_array outTexture [[texture(1)]], + const device float4 *new_scale [[buffer(0)]], + const device float4 *new_biase [[buffer(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; + } + + float4 input; + float4 output; + constexpr sampler sample(coord::pixel, filter::nearest, address::clamp_to_zero); + input = inTexture.sample(sample, gid.x, gid.y, gid.z); + output = fmax(input * new_scale[gid.z] + new_biase[gid.z], 0.0); + outTexture.write(output, gid.xy, gid.z); + +} diff --git a/metal/paddle-mobile/paddle-mobile/PaddleMobile.swift b/metal/paddle-mobile/paddle-mobile/PaddleMobile.swift new file mode 100644 index 0000000000000000000000000000000000000000..8e28fb2c2d77c4324cb38bbc9c0e178cc1780697 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/PaddleMobile.swift @@ -0,0 +1,175 @@ +// +// PaddleMobile.swift +// paddle-mobile-demo +// +// Created by liuRuiLong on 2018/9/5. +// Copyright © 2018年 orange. All rights reserved. +// + +import Metal +import MetalKit +import Foundation + +public enum Platform{ + case CPU, GPU +} + +class ScaleKernel: CusomKernel { + init(device: MTLDevice, shape: Shape) { + super.init(device: device, inFunctionName: "scale", outputDim: shape, usePaddleMobileLib: false) + } +} + +public protocol Net { + var except: Int { get } + var dim: (n: Int, h: Int, w: Int, c: Int) { get } + var preprocessKernel: CusomKernel { get } +// var paramPointer: UnsafeMutableRawPointer { get } +// var paramSize: Int { get } +// var modelPointer: UnsafeMutableRawPointer { get } +// var modelSize: Int { get } + var modelPath: String { get } + var paramPath: String { get } + var modelDir: String { get } + func resultStr(res: [Float]) -> String + func fetchResult(paddleMobileRes: ResultHolder) -> [Float32] +} + +extension Net { + func fetchResult(paddleMobileRes: ResultHolder) -> [Float32] { + return paddleMobileRes.resultArr + } +} + +public class Runner { + var program: Program? + var executor: Executor? + var queue: MTLCommandQueue? + var textureLoader: MTKTextureLoader? + let net: Net + let device: MTLDevice? + let platform: Platform + /** + * inNet: 需要运行的网络 + * commandQueue: GPU 是需要传入 + * inPlatform: 需要使用的平台, GPU or CPU + */ + public init(inNet: Net, commandQueue: MTLCommandQueue?, inPlatform: Platform) { + net = inNet + queue = commandQueue + device = queue?.device + platform = inPlatform + if let inDevice = device { + textureLoader = MTKTextureLoader.init(device: inDevice) + } + } + + /** + * load 模型, 返回 true 可进行预测 + */ + public func load() -> Bool { + if platform == .GPU { + guard let inDevice = device, let inQueue = queue else { + print(" paddle mobile gpu load error, need MTLCommandQueue") + return false + } + let loader = Loader.init() + do { + program = try loader.load(device: inDevice, modelPath: net.modelPath, paraPath: net.paramPath) + executor = try Executor.init(inDevice: inDevice, inQueue: inQueue, inProgram: program!) + } catch let error { + print(error) + return false + } + } else { + print(" need implementation ") + return false + } + return true + } + + /** + * CPU GPU 通用版本 predict + * cgImage: 需要预测的图片 + * ( _ success: Bool, _ time:TimeInterval, _ resultArray: [Float32]) -> Void : 回调闭包, 三个参数分别为: 是否成功, 预测耗时, 结果数组 + */ + public func predict(cgImage: CGImage, completion: @escaping ( _ success: Bool, _ time:TimeInterval, _ resultArray: [Float32]) -> Void) { + if platform == .GPU { + getTexture(image: cgImage) { [weak self] (texture) in + guard let SSelf = self else { + fatalError() + } + SSelf.predict(texture: texture, completion: completion) + } + } else if platform == .CPU { + + } + } + /** + * GPU 版本 predict + * texture: 需要预测的 texture 需要做过预处理 + * ( _ success: Bool, _ time:TimeInterval, _ resultArray: [Float32]) -> Void : 回调闭包, 三个参数分别为: 是否成功, 预测耗时, 结果数组 + */ + public func predict(texture: MTLTexture, completion: @escaping ( _ success: Bool, _ time:TimeInterval, _ resultArray: [Float32]) -> Void) { + do { + try self.executor?.predict(input: texture, dim: [self.net.dim.n, self.net.dim.h, self.net.dim.w, self.net.dim.c], completionHandle: { [weak self] (res) in + guard let SSelf = self else { + fatalError( " self nil " ) + } + let resultArray = SSelf.net.fetchResult(paddleMobileRes: res) + completion(true, res.elapsedTime, resultArray) + }, preProcessKernle: self.net.preprocessKernel, except: self.net.except) + } catch let error { + print(error) + completion(false, 0.0, []) + return + } + } + /* + * 清理内存, 调用此函数后, 不能再使用, 需重新 load + */ + public func clear() { + executor?.clear() + executor = nil + program = nil + } + + /* + * 获取 texture, 对 texture 进行预处理, GPU 预测时使用 + */ + public func getTexture(image: CGImage, getTexture: @escaping (MTLTexture) -> Void) { + let texture = try? textureLoader?.newTexture(cgImage: image, options: [:]) ?! " texture loader error" + scaleTexture(input: texture!, size: (net.dim.w, net.dim.h), complete: getTexture) + } + + func scaleTexture(input: MTLTexture, size:(width: Int, height: Int), complete: @escaping (MTLTexture) -> Void) { + + guard let inQueue = queue, let inDevice = device else { + fatalError( " queue or devcie nil " ) + } + + guard let buffer = inQueue.makeCommandBuffer() else { + fatalError( " make buffer error" ) + } + + let scaleKernel = ScaleKernel.init(device: inDevice, shape: CusomKernel.Shape.init(inWidth: size.width, inHeight: size.height, inChannel: 3)) + + do { + try scaleKernel.compute(inputTexuture: input, commandBuffer: buffer) + } catch let error { + print(error) + fatalError() + } + + buffer.addCompletedHandler { (buffer) in + complete(scaleKernel.outputTexture) + } + buffer.commit() + } +} + + + + + + diff --git a/metal/paddle-mobile/paddle-mobile/Executor.swift b/metal/paddle-mobile/paddle-mobile/framework/Executor.swift similarity index 100% rename from metal/paddle-mobile/paddle-mobile/Executor.swift rename to metal/paddle-mobile/paddle-mobile/framework/Executor.swift diff --git a/metal/paddle-mobile/paddle-mobile/Loader.swift b/metal/paddle-mobile/paddle-mobile/framework/Loader.swift similarity index 100% rename from metal/paddle-mobile/paddle-mobile/Loader.swift rename to metal/paddle-mobile/paddle-mobile/framework/Loader.swift