提交 e273b9dd 编写于 作者: D dolphin8

finish

...@@ -14,6 +14,9 @@ ...@@ -14,6 +14,9 @@
FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8520E11C550081E9F8 /* Main.storyboard */; }; FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8520E11C550081E9F8 /* Main.storyboard */; };
FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8820E11C560081E9F8 /* Assets.xcassets */; }; FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8820E11C560081E9F8 /* Assets.xcassets */; };
FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */; }; FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */; };
FC27991021341CE5000B6BAD /* Net.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC27990F21341CE5000B6BAD /* Net.swift */; };
FC3C800F2133F46600D1295E /* MobileNetSSD.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3C800E2133F46600D1295E /* MobileNetSSD.swift */; };
FC3C80112133F4AB00D1295E /* MobileNet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3C80102133F4AB00D1295E /* MobileNet.swift */; };
FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */ = {isa = PBXBuildFile; fileRef = FC918190211DBC3500B6F354 /* paddle-mobile.png */; }; FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */ = {isa = PBXBuildFile; fileRef = FC918190211DBC3500B6F354 /* paddle-mobile.png */; };
FC918193211DC70500B6F354 /* iphone.JPG in Resources */ = {isa = PBXBuildFile; fileRef = FC918192211DC70500B6F354 /* iphone.JPG */; }; FC918193211DC70500B6F354 /* iphone.JPG in Resources */ = {isa = PBXBuildFile; fileRef = FC918192211DC70500B6F354 /* iphone.JPG */; };
FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FCA3A16021313E1F00084FE5 /* hand.jpg */; }; FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FCA3A16021313E1F00084FE5 /* hand.jpg */; };
...@@ -22,7 +25,6 @@ ...@@ -22,7 +25,6 @@
FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC542122EF5400D94F7E /* MetalHelper.swift */; }; FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC542122EF5400D94F7E /* MetalHelper.swift */; };
FCD04E6320F3146B0007374F /* params in Resources */ = {isa = PBXBuildFile; fileRef = FCD04E6120F3146A0007374F /* params */; }; FCD04E6320F3146B0007374F /* params in Resources */ = {isa = PBXBuildFile; fileRef = FCD04E6120F3146A0007374F /* params */; };
FCD04E6420F3146B0007374F /* model in Resources */ = {isa = PBXBuildFile; fileRef = FCD04E6220F3146A0007374F /* model */; }; FCD04E6420F3146B0007374F /* model in Resources */ = {isa = PBXBuildFile; fileRef = FCD04E6220F3146A0007374F /* model */; };
FCDFD3FB211D72C3005AB38B /* ModelHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCDFD3FA211D72C3005AB38B /* ModelHelper.swift */; };
FCDFD41B211D91C7005AB38B /* synset.txt in Resources */ = {isa = PBXBuildFile; fileRef = FCDFD41A211D91C7005AB38B /* synset.txt */; }; FCDFD41B211D91C7005AB38B /* synset.txt in Resources */ = {isa = PBXBuildFile; fileRef = FCDFD41A211D91C7005AB38B /* synset.txt */; };
FCEBEC2C20E1391F00C0B14D /* paddle_mobile.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */; }; FCEBEC2C20E1391F00C0B14D /* paddle_mobile.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */; };
FCEBEC2D20E1391F00C0B14D /* paddle_mobile.framework in Embed Frameworks */ = {isa = PBXBuildFile; fileRef = FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */; settings = {ATTRIBUTES = (CodeSignOnCopy, RemoveHeadersOnCopy, ); }; }; FCEBEC2D20E1391F00C0B14D /* paddle_mobile.framework in Embed Frameworks */ = {isa = PBXBuildFile; fileRef = FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */; settings = {ATTRIBUTES = (CodeSignOnCopy, RemoveHeadersOnCopy, ); }; };
...@@ -55,6 +57,9 @@ ...@@ -55,6 +57,9 @@
FC039B8820E11C560081E9F8 /* Assets.xcassets */ = {isa = PBXFileReference; lastKnownFileType = folder.assetcatalog; path = Assets.xcassets; sourceTree = "<group>"; }; FC039B8820E11C560081E9F8 /* Assets.xcassets */ = {isa = PBXFileReference; lastKnownFileType = folder.assetcatalog; path = Assets.xcassets; sourceTree = "<group>"; };
FC039B8B20E11C560081E9F8 /* Base */ = {isa = PBXFileReference; lastKnownFileType = file.storyboard; name = Base; path = Base.lproj/LaunchScreen.storyboard; sourceTree = "<group>"; }; FC039B8B20E11C560081E9F8 /* Base */ = {isa = PBXFileReference; lastKnownFileType = file.storyboard; name = Base; path = Base.lproj/LaunchScreen.storyboard; sourceTree = "<group>"; };
FC039B8D20E11C560081E9F8 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; sourceTree = "<group>"; }; FC039B8D20E11C560081E9F8 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; sourceTree = "<group>"; };
FC27990F21341CE5000B6BAD /* Net.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Net.swift; sourceTree = "<group>"; };
FC3C800E2133F46600D1295E /* MobileNetSSD.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNetSSD.swift; sourceTree = "<group>"; };
FC3C80102133F4AB00D1295E /* MobileNet.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNet.swift; sourceTree = "<group>"; };
FC918190211DBC3500B6F354 /* paddle-mobile.png */ = {isa = PBXFileReference; lastKnownFileType = image.png; path = "paddle-mobile.png"; sourceTree = "<group>"; }; FC918190211DBC3500B6F354 /* paddle-mobile.png */ = {isa = PBXFileReference; lastKnownFileType = image.png; path = "paddle-mobile.png"; sourceTree = "<group>"; };
FC918192211DC70500B6F354 /* iphone.JPG */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = iphone.JPG; sourceTree = "<group>"; }; FC918192211DC70500B6F354 /* iphone.JPG */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = iphone.JPG; sourceTree = "<group>"; };
FCA3A16021313E1F00084FE5 /* hand.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = hand.jpg; sourceTree = "<group>"; }; FCA3A16021313E1F00084FE5 /* hand.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = hand.jpg; sourceTree = "<group>"; };
...@@ -63,7 +68,6 @@ ...@@ -63,7 +68,6 @@
FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = "<group>"; }; FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = "<group>"; };
FCD04E6120F3146A0007374F /* params */ = {isa = PBXFileReference; lastKnownFileType = file; path = params; sourceTree = "<group>"; }; FCD04E6120F3146A0007374F /* params */ = {isa = PBXFileReference; lastKnownFileType = file; path = params; sourceTree = "<group>"; };
FCD04E6220F3146A0007374F /* model */ = {isa = PBXFileReference; lastKnownFileType = file; path = model; sourceTree = "<group>"; }; FCD04E6220F3146A0007374F /* model */ = {isa = PBXFileReference; lastKnownFileType = file; path = model; sourceTree = "<group>"; };
FCDFD3FA211D72C3005AB38B /* ModelHelper.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ModelHelper.swift; sourceTree = "<group>"; };
FCDFD41A211D91C7005AB38B /* synset.txt */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = synset.txt; sourceTree = "<group>"; }; FCDFD41A211D91C7005AB38B /* synset.txt */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = synset.txt; sourceTree = "<group>"; };
FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; path = paddle_mobile.framework; sourceTree = BUILT_PRODUCTS_DIR; }; FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; path = paddle_mobile.framework; sourceTree = BUILT_PRODUCTS_DIR; };
FCEEE7D3210627A000444BEC /* banana.jpeg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = banana.jpeg; sourceTree = "<group>"; }; FCEEE7D3210627A000444BEC /* banana.jpeg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = banana.jpeg; sourceTree = "<group>"; };
...@@ -131,7 +135,9 @@ ...@@ -131,7 +135,9 @@
FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */, FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */,
FC039B8D20E11C560081E9F8 /* Info.plist */, FC039B8D20E11C560081E9F8 /* Info.plist */,
FCBCCC542122EF5400D94F7E /* MetalHelper.swift */, FCBCCC542122EF5400D94F7E /* MetalHelper.swift */,
FCDFD3FA211D72C3005AB38B /* ModelHelper.swift */, FC3C800E2133F46600D1295E /* MobileNetSSD.swift */,
FC3C80102133F4AB00D1295E /* MobileNet.swift */,
FC27990F21341CE5000B6BAD /* Net.swift */,
); );
path = "paddle-mobile-demo"; path = "paddle-mobile-demo";
sourceTree = "<group>"; sourceTree = "<group>";
...@@ -300,9 +306,11 @@ ...@@ -300,9 +306,11 @@
buildActionMask = 2147483647; buildActionMask = 2147483647;
files = ( files = (
FC039B8420E11C550081E9F8 /* ViewController.swift in Sources */, FC039B8420E11C550081E9F8 /* ViewController.swift in Sources */,
FCDFD3FB211D72C3005AB38B /* ModelHelper.swift in Sources */,
FC013928210204A3008100E3 /* PreProcessKernel.metal in Sources */, FC013928210204A3008100E3 /* PreProcessKernel.metal in Sources */,
FC27991021341CE5000B6BAD /* Net.swift in Sources */,
FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */, FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */,
FC3C80112133F4AB00D1295E /* MobileNet.swift in Sources */,
FC3C800F2133F46600D1295E /* MobileNetSSD.swift in Sources */,
FC039B8220E11C550081E9F8 /* AppDelegate.swift in Sources */, FC039B8220E11C550081E9F8 /* AppDelegate.swift in Sources */,
); );
runOnlyForDeploymentPostprocessing = 0; runOnlyForDeploymentPostprocessing = 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 paddle_mobile
class MobileNet: Net{
var program: Program?
var executor: Executor<Float32>?
let except: Int = 0
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) {
if let filePath = inBundle.path(forResource: fileName, ofType: type) {
let string = try! String.init(contentsOfFile: filePath)
contents = string.components(separatedBy: CharacterSet.newlines).filter{$0.count > 10}.map{
String($0[$0.index($0.startIndex, offsetBy: 10)...])
}
}else{
fatalError("no file call \(fileName)")
}
}
subscript(index: Int) -> String {
return contents[index]
}
}
let labels = PreWords.init(fileName: "synset")
func resultStr(res: [Float]) -> String {
var s: [String] = []
res.top(r: 5).enumerated().forEach{
s.append(String(format: "%d: %@ (%3.2f%%)", $0 + 1, labels[$1.0], $1.1 * 100))
}
return s.joined(separator: "\n")
}
var preprocessKernel: CusomKernel
let dim = [1, 224, 224, 3]
let modelPath: String
let paramPath: String
let modelDir: String
init() {
modelPath = Bundle.main.path(forResource: "model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null"
modelDir = ""
preprocessKernel = MobilenetPreProccess.init(device: MetalHelper.shared.device)
}
}
// /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
// ModelHelper.swift
// paddle-mobile-demo Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// Created by liuRuiLong on 2018/8/10. You may obtain a copy of the License at
// Copyright © 2018年 orange. All rights reserved.
// 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 UIKit
import MetalKit
import Foundation import Foundation
import paddle_mobile import paddle_mobile
import MetalPerformanceShaders
let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()] class MobileNet_ssd_hand: Net{
enum SupportModel: String{
case mobilenet = "mobilenet"
case mobilenet_ssd = "mobilenetssd"
static func supportedModels() -> [SupportModel] {
return [.mobilenet, .mobilenet_ssd]
}
}
protocol Net {
var dim: [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>) -> [Float32]
}
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
getTexture(resTexture)
}
}
func fetchResult(paddleMobileRes: ResultHolder<Float32>) -> [Float32] { var program: Program?
return paddleMobileRes.resultArr
}
}
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 executor: Executor<Float32>?
var contents: [String] = []
init(fileName: String, type: String = "txt", inBundle: Bundle = Bundle.main) {
if let filePath = inBundle.path(forResource: fileName, ofType: type) {
let string = try! String.init(contentsOfFile: filePath)
contents = string.components(separatedBy: CharacterSet.newlines).filter{$0.count > 10}.map{
String($0[$0.index($0.startIndex, offsetBy: 10)...])
}
}else{
fatalError("no file call \(fileName)")
}
}
subscript(index: Int) -> String {
return contents[index]
}
}
let labels = PreWords.init(fileName: "synset")
func resultStr(res: [Float]) -> String {
var s: [String] = []
res.top(r: 5).enumerated().forEach{
s.append(String(format: "%d: %@ (%3.2f%%)", $0 + 1, labels[$1.0], $1.1 * 100))
}
return s.joined(separator: "\n")
}
var preprocessKernel: CusomKernel
let dim = [1, 224, 224, 3]
let modelPath: String
let paramPath: String
let modelDir: String
init() { let except: Int = 2
modelPath = Bundle.main.path(forResource: "model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null"
modelDir = ""
preprocessKernel = MobilenetPreProccess.init(device: MetalHelper.shared.device)
}
}
struct MobileNet_ssd_hand: Net{
class MobilenetssdPreProccess: CusomKernel { class MobilenetssdPreProccess: CusomKernel {
init(device: MTLDevice) { init(device: MTLDevice) {
let s = CusomKernel.Shape.init(inWidth: 300, inHeight: 300, inChannel: 3) let s = CusomKernel.Shape.init(inWidth: 300, inHeight: 300, inChannel: 3)
...@@ -105,7 +31,6 @@ struct MobileNet_ssd_hand: Net{ ...@@ -105,7 +31,6 @@ struct MobileNet_ssd_hand: Net{
func resultStr(res: [Float]) -> String { func resultStr(res: [Float]) -> String {
return "哈哈哈, 还没好" return "哈哈哈, 还没好"
// fatalError()
} }
func bboxArea(box: [Float32], normalized: Bool) -> Float32 { func bboxArea(box: [Float32], normalized: Bool) -> Float32 {
...@@ -141,10 +66,18 @@ struct MobileNet_ssd_hand: Net{ ...@@ -141,10 +66,18 @@ struct MobileNet_ssd_hand: Net{
} }
func fetchResult(paddleMobileRes: ResultHolder<Float32>) -> [Float32]{ func fetchResult(paddleMobileRes: ResultHolder<Float32>) -> [Float32]{
let scores = paddleMobileRes.intermediateResults![0] as! Texture<Float32> guard let interRes = paddleMobileRes.intermediateResults else {
let bbox = paddleMobileRes.intermediateResults![1] as! Texture<Float32> fatalError(" need have inter result ")
// let bbox = paddleMobileRes["box_coder_0.tmp_0"] ?! " no bbox " }
// let scores = paddleMobileRes["transpose_12.tmp_0"] ?! " no scores "
guard let scores = interRes["Scores"], scores.count > 0, let score = scores[0] as? Texture<Float32> else {
fatalError(" need score ")
}
guard let bboxs = interRes["BBoxes"], bboxs.count > 0, let bbox = bboxs[0] as? Texture<Float32> else {
fatalError()
}
let score_thredshold: Float32 = 0.01 let score_thredshold: Float32 = 0.01
let nms_top_k = 400 let nms_top_k = 400
let keep_top_k = 200 let keep_top_k = 200
...@@ -155,43 +88,14 @@ struct MobileNet_ssd_hand: Net{ ...@@ -155,43 +88,14 @@ struct MobileNet_ssd_hand: Net{
return f return f
} }
let scoresArr = scores.metalTexture.floatArray { (f) -> Float32 in let scoreFormatArr: [Float32] = score.metalTexture.realNHWC(dim: (n: score.originDim[0], h: score.originDim[1], w: score.originDim[2], c: score.originDim[3]))
return f
}
var scoreFormatArr: [Float32] = []
var outputArr: [Float32] = [] var outputArr: [Float32] = []
let cNumOfOneClass = score.tensorDim[2] // 1917
let numOfOneC = (scores.tensorDim[2] + 3) / 4 // 480
let cNumOfOneClass = scores.tensorDim[2] // 1917
let cPaddedNumOfOneClass = numOfOneC * 4 // 1920
let boxSize = bbox.tensorDim[2] // 4 let boxSize = bbox.tensorDim[2] // 4
let classNum = scores.tensorDim[1] // 7 let classNum = score.tensorDim[1] // 7
let classNumOneTexture = classNum * 4 // 28
for c in 0..<classNum {
for n in 0..<numOfOneC {
let to = n * classNumOneTexture + c * 4
if n == numOfOneC - 1 {
for i in 0..<(4 - (cPaddedNumOfOneClass - cNumOfOneClass)) {
scoreFormatArr.append(scoresArr[to + i])
}
} else {
scoreFormatArr.append(scoresArr[to])
scoreFormatArr.append(scoresArr[to + 1])
scoreFormatArr.append(scoresArr[to + 2])
scoreFormatArr.append(scoresArr[to + 3])
}
}
}
var selectedIndexs: [Int : [(Int, Float32)]] = [:] var selectedIndexs: [Int : [(Int, Float32)]] = [:]
var numDet: Int = 0 var numDet: Int = 0
for i in 0..<classNum { for i in 0..<classNum {
var sliceScore = Array<Float32>(scoreFormatArr[(i * cNumOfOneClass)..<((i + 1) * cNumOfOneClass)]) var sliceScore = Array<Float32>(scoreFormatArr[(i * cNumOfOneClass)..<((i + 1) * cNumOfOneClass)])
...@@ -249,7 +153,7 @@ struct MobileNet_ssd_hand: Net{ ...@@ -249,7 +153,7 @@ struct MobileNet_ssd_hand: Net{
} }
scoreIndexPairs.sort { $0.0 > $1.0 } scoreIndexPairs.sort { $0.0 > $1.0 }
if scoreIndexPairs.count > keep_top_k { if scoreIndexPairs.count > keep_top_k {
scoreIndexPairs.removeLast(scoreIndexPairs.count - keep_top_k) scoreIndexPairs.removeLast(scoreIndexPairs.count - keep_top_k)
} }
...@@ -259,9 +163,9 @@ struct MobileNet_ssd_hand: Net{ ...@@ -259,9 +163,9 @@ struct MobileNet_ssd_hand: Net{
// label: scoreIndexPair.1.0 // label: scoreIndexPair.1.0
let label = scoreIndexPair.1.0 let label = scoreIndexPair.1.0
if newIndices[label] != nil { if newIndices[label] != nil {
newIndices[label]?.append((scoreIndexPair.1.0, scoreIndexPair.0)) newIndices[label]?.append((scoreIndexPair.1.1, scoreIndexPair.0))
} else { } else {
newIndices[label] = [(scoreIndexPair.1.0, scoreIndexPair.0)] newIndices[label] = [(scoreIndexPair.1.1, scoreIndexPair.0)]
} }
} }
...@@ -275,7 +179,6 @@ struct MobileNet_ssd_hand: Net{ ...@@ -275,7 +179,6 @@ struct MobileNet_ssd_hand: Net{
} }
} }
print(" fuck success !") print(" fuck success !")
print(outputArr)
return outputArr return outputArr
} }
...@@ -292,4 +195,3 @@ struct MobileNet_ssd_hand: Net{ ...@@ -292,4 +195,3 @@ struct MobileNet_ssd_hand: Net{
preprocessKernel = MobilenetssdPreProccess.init(device: MetalHelper.shared.device) preprocessKernel = MobilenetssdPreProccess.init(device: MetalHelper.shared.device)
} }
} }
//
// Net.swift
// paddle-mobile-demo
//
// Created by liuRuiLong on 2018/8/27.
// Copyright © 2018年 orange. All rights reserved.
//
import Foundation
import UIKit
import MetalKit
import Foundation
import paddle_mobile
import MetalPerformanceShaders
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, .mobilenet_ssd]
}
}
protocol Net {
var program: Program? { get set }
var executor: Executor<Float32>? { get set }
var except: Int { get }
var dim: [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>) -> [Float32]
mutating func load() throws
func predict(inTexture: MTLTexture, completion: @escaping ([Float32]) -> Void) throws
mutating func clear()
}
extension Net {
mutating func load() throws {
let queue = MetalHelper.shared.queue
let loader = Loader<Float32>.init()
do {
program = try loader.load(device: MetalHelper.shared.device, modelPath: modelPath, paraPath: paramPath)
executor = try Executor<Float32>.init(inDevice: MetalHelper.shared.device, inQueue: queue, inProgram: program!)
} catch let error {
throw error
}
}
func predict(inTexture: MTLTexture, completion: @escaping ([Float32]) -> Void) throws {
guard let inExecutor = executor else {
fatalError(" 请先 load ")
}
try inExecutor.predict(input: inTexture, dim: dim, completionHandle: { (result) in
let resultArr = self.fetchResult(paddleMobileRes: result)
completion(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: (224, 224)) { (resTexture) in
getTexture(resTexture)
}
}
func fetchResult(paddleMobileRes: ResultHolder<Float32>) -> [Float32] {
return paddleMobileRes.resultArr
}
// func predict()
}
...@@ -26,25 +26,21 @@ class ViewController: UIViewController { ...@@ -26,25 +26,21 @@ class ViewController: UIViewController {
@IBOutlet weak var modelPickerView: UIPickerView! @IBOutlet weak var modelPickerView: UIPickerView!
@IBOutlet weak var threadPickerView: UIPickerView! @IBOutlet weak var threadPickerView: UIPickerView!
var selectImage: UIImage? var selectImage: UIImage?
var program: Program?
var executor: Executor<Float32>?
var modelType: SupportModel = SupportModel.supportedModels()[0] var modelType: SupportModel = SupportModel.supportedModels()[0]
var toPredictTexture: MTLTexture? var toPredictTexture: MTLTexture?
var modelHelper: Net { var net: Net {
return modelHelperMap[modelType] ?! " has no this type " get {
return modelHelperMap[modelType] ?! " has no this type "
}
set {
}
} }
var threadNum = 1 var threadNum = 1
@IBAction func loadAct(_ sender: Any) { @IBAction func loadAct(_ sender: Any) {
let inModelHelper = modelHelper
let queue = MetalHelper.shared.queue
let loader = Loader<Float32>.init()
do { do {
let modelPath = inModelHelper.modelPath try self.net.load()
let paraPath = inModelHelper.paramPath
program = try loader.load(device: MetalHelper.shared.device, modelPath: modelPath, paraPath: paraPath)
executor = try Executor<Float32>.init(inDevice: MetalHelper.shared.device, inQueue: queue, inProgram: program!)
} catch let error { } catch let error {
print(error) print(error)
} }
...@@ -58,9 +54,7 @@ class ViewController: UIViewController { ...@@ -58,9 +54,7 @@ class ViewController: UIViewController {
} }
@IBAction func clearAct(_ sender: Any) { @IBAction func clearAct(_ sender: Any) {
executor?.clear() net.clear()
program = nil
executor = nil
} }
@IBAction func predictAct(_ sender: Any) { @IBAction func predictAct(_ sender: Any) {
...@@ -68,39 +62,54 @@ class ViewController: UIViewController { ...@@ -68,39 +62,54 @@ class ViewController: UIViewController {
resultTextView.text = "请选择图片 ! " resultTextView.text = "请选择图片 ! "
return return
} }
guard let inExecutor = executor else {
resultTextView.text = "请先 load ! "
return
}
do { do {
let max = 1 try net.predict(inTexture: inTexture) { [weak self] (result) in
var startDate = Date.init() guard let sSelf = self else {
for i in 0..<max { fatalError()
try inExecutor.predict(input: inTexture, expect: modelHelper.dim, completionHandle: { [weak self] (result) in }
guard let sSelf = self else { print(result)
fatalError() let resultStr = sSelf.net.resultStr(res: result)
} DispatchQueue.main.async {
sSelf.resultTextView.text = resultStr
if i == (max / 2 - 1) { }
startDate = Date.init()
}
let resultArr = sSelf.modelHelper.fetchResult(paddleMobileRes: result)
if i == max - 1 {
let time = Date.init().timeIntervalSince(startDate)
DispatchQueue.main.async {
sSelf.resultTextView.text = sSelf.modelHelper.resultStr(res: resultArr)
sSelf.elapsedTimeLabel.text = "平均耗时: \(time/Double(max/2) * 1000.0) ms"
}
}
}, preProcessKernle: self.modelHelper.preprocessKernel, except: 2)
} }
} catch let error { } catch let error {
print(error) print(error)
} }
// guard let inExecutor = executor else {
// resultTextView.text = "请先 load ! "
// return
// }
//
// do {
// let max = 1
// var startDate = Date.init()
// for i in 0..<max {
// try inExecutor.predict(input: inTexture, dim: modelHelper.dim, completionHandle: { [weak self] (result) in
// guard let sSelf = self else {
// fatalError()
// }
//
// if i == (max / 2 - 1) {
// startDate = Date.init()
// }
//
// let resultArr = sSelf.modelHelper.fetchResult(paddleMobileRes: result)
//
// if i == max - 1 {
// let time = Date.init().timeIntervalSince(startDate)
// DispatchQueue.main.async {
// sSelf.resultTextView.text = sSelf.modelHelper.resultStr(res: resultArr)
// sSelf.elapsedTimeLabel.text = "平均耗时: \(time/Double(max/2) * 1000.0) ms"
// }
// }
// }, preProcessKernle: self.modelHelper.preprocessKernel, except: modelHelper.except)
// }
// } catch let error {
// print(error)
// }
} }
override func viewDidLoad() { override func viewDidLoad() {
...@@ -112,7 +121,7 @@ class ViewController: UIViewController { ...@@ -112,7 +121,7 @@ class ViewController: UIViewController {
selectImage = UIImage.init(named: "hand.jpg") selectImage = UIImage.init(named: "hand.jpg")
selectImageView.image = selectImage selectImageView.image = selectImage
modelHelper.getTexture(image: selectImage!.cgImage!) {[weak self] (texture) in net.getTexture(image: selectImage!.cgImage!) {[weak self] (texture) in
self?.toPredictTexture = texture self?.toPredictTexture = texture
} }
} }
...@@ -168,7 +177,7 @@ extension ViewController: UIImagePickerControllerDelegate, UINavigationControll ...@@ -168,7 +177,7 @@ extension ViewController: UIImagePickerControllerDelegate, UINavigationControll
} }
sSelf.selectImage = image sSelf.selectImage = image
sSelf.selectImageView.image = image sSelf.selectImageView.image = image
sSelf.modelHelper.getTexture(image: image.cgImage!, getTexture: { (texture) in sSelf.net.getTexture(image: image.cgImage!, getTexture: { (texture) in
sSelf.toPredictTexture = texture sSelf.toPredictTexture = texture
}) })
} }
......
...@@ -38,6 +38,7 @@ ...@@ -38,6 +38,7 @@
FC0E2DC020EE461F009C1FAC /* ElementwiseAddKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */; }; FC0E2DC020EE461F009C1FAC /* ElementwiseAddKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */; };
FC1B16B320EC9A4F00678B91 /* Kernels.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC1B16B220EC9A4F00678B91 /* Kernels.metal */; }; FC1B16B320EC9A4F00678B91 /* Kernels.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC1B16B220EC9A4F00678B91 /* Kernels.metal */; };
FC1B186620ECF1C600678B91 /* ResizeKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC1B186520ECF1C600678B91 /* ResizeKernel.swift */; }; FC1B186620ECF1C600678B91 /* ResizeKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC1B186520ECF1C600678B91 /* ResizeKernel.swift */; };
FC27990E21341016000B6BAD /* BoxCoder.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC27990D21341016000B6BAD /* BoxCoder.metal */; };
FC3602CC2108819F00FACB58 /* PaddleMobileUnitTest.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */; }; FC3602CC2108819F00FACB58 /* PaddleMobileUnitTest.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */; };
FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74820F0B954007C0C6D /* ConvKernel.metal */; }; FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74820F0B954007C0C6D /* ConvKernel.metal */; };
FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */; }; FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */; };
...@@ -122,6 +123,7 @@ ...@@ -122,6 +123,7 @@
FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ElementwiseAddKernel.swift; sourceTree = "<group>"; }; FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ElementwiseAddKernel.swift; sourceTree = "<group>"; };
FC1B16B220EC9A4F00678B91 /* Kernels.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = Kernels.metal; sourceTree = "<group>"; }; FC1B16B220EC9A4F00678B91 /* Kernels.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = Kernels.metal; sourceTree = "<group>"; };
FC1B186520ECF1C600678B91 /* ResizeKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ResizeKernel.swift; sourceTree = "<group>"; }; FC1B186520ECF1C600678B91 /* ResizeKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ResizeKernel.swift; sourceTree = "<group>"; };
FC27990D21341016000B6BAD /* BoxCoder.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BoxCoder.metal; sourceTree = "<group>"; };
FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PaddleMobileUnitTest.swift; sourceTree = "<group>"; }; FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PaddleMobileUnitTest.swift; sourceTree = "<group>"; };
FC4CB74820F0B954007C0C6D /* ConvKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvKernel.metal; sourceTree = "<group>"; }; FC4CB74820F0B954007C0C6D /* ConvKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvKernel.metal; sourceTree = "<group>"; };
FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ProgramOptimize.swift; sourceTree = "<group>"; }; FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ProgramOptimize.swift; sourceTree = "<group>"; };
...@@ -351,6 +353,7 @@ ...@@ -351,6 +353,7 @@
FCEB6837212F00B100D2448E /* metal */ = { FCEB6837212F00B100D2448E /* metal */ = {
isa = PBXGroup; isa = PBXGroup;
children = ( children = (
FC27990D21341016000B6BAD /* BoxCoder.metal */,
FC1B16B220EC9A4F00678B91 /* Kernels.metal */, FC1B16B220EC9A4F00678B91 /* Kernels.metal */,
FC4CB74820F0B954007C0C6D /* ConvKernel.metal */, FC4CB74820F0B954007C0C6D /* ConvKernel.metal */,
4AF928762133F1DB005B6C3A /* BoxCoder.metal */, 4AF928762133F1DB005B6C3A /* BoxCoder.metal */,
......
...@@ -442,10 +442,5 @@ public extension MTLBuffer { ...@@ -442,10 +442,5 @@ public extension MTLBuffer {
} }
return array; return array;
} }
} }
...@@ -52,9 +52,6 @@ extension Float16: PrecisionType { ...@@ -52,9 +52,6 @@ extension Float16: PrecisionType {
public init(inFloat: Float32) { public init(inFloat: Float32) {
self = Int16(inFloat) self = Int16(inFloat)
} }
} }
extension Float32: PrecisionType { extension Float32: PrecisionType {
...@@ -214,10 +211,6 @@ extension DataLayout: Equatable { ...@@ -214,10 +211,6 @@ extension DataLayout: Equatable {
} }
} }
public protocol Variant: CustomStringConvertible, CustomDebugStringConvertible { public protocol Variant: CustomStringConvertible, CustomDebugStringConvertible {
} }
......
...@@ -14,14 +14,14 @@ ...@@ -14,14 +14,14 @@
import Foundation import Foundation
let testTo = 94 let testTo = 99
public class ResultHolder<P: PrecisionType> { public class ResultHolder<P: PrecisionType> {
public let dim: [Int] public let dim: [Int]
public let resultArr: [P] public let resultArr: [P]
public var intermediateResults: [Variant]? public var intermediateResults: [String : [Variant]]?
public let elapsedTime: Double public let elapsedTime: Double
public init(inDim: [Int], inResult: [P], inElapsedTime: Double, inIntermediateResults: [Variant]? = nil) { public init(inDim: [Int], inResult: [P], inElapsedTime: Double, inIntermediateResults: [String : [Variant]]? = nil) {
dim = inDim dim = inDim
resultArr = inResult resultArr = inResult
elapsedTime = inElapsedTime elapsedTime = inElapsedTime
...@@ -62,7 +62,7 @@ public class Executor<P: PrecisionType> { ...@@ -62,7 +62,7 @@ public class Executor<P: PrecisionType> {
queue = inQueue queue = inQueue
for block in inProgram.programDesc.blocks { for block in inProgram.programDesc.blocks {
//block.ops.count //block.ops.count
for i in 0..<(testTo + 1) { for i in 0..<block.ops.count {
let op = block.ops[i] let op = block.ops[i]
do { do {
let op = try OpCreator<P>.shared.creat(device: inDevice, opDesc: op, scope: inProgram.scope) let op = try OpCreator<P>.shared.creat(device: inDevice, opDesc: op, scope: inProgram.scope)
...@@ -75,7 +75,7 @@ public class Executor<P: PrecisionType> { ...@@ -75,7 +75,7 @@ public class Executor<P: PrecisionType> {
} }
} }
public func predict(input: MTLTexture, expect: [Int], completionHandle: @escaping (ResultHolder<P>) -> Void, preProcessKernle: CusomKernel? = nil, except: Int = 0) throws { public func predict(input: MTLTexture, dim: [Int], completionHandle: @escaping (ResultHolder<P>) -> Void, preProcessKernle: CusomKernel? = nil, except: Int = 0) throws {
guard let buffer = queue.makeCommandBuffer() else { guard let buffer = queue.makeCommandBuffer() else {
throw PaddleMobileError.predictError(message: "CommandBuffer is nil") throw PaddleMobileError.predictError(message: "CommandBuffer is nil")
} }
...@@ -92,7 +92,7 @@ public class Executor<P: PrecisionType> { ...@@ -92,7 +92,7 @@ public class Executor<P: PrecisionType> {
} }
let beforeDate = Date.init() let beforeDate = Date.init()
let inputTexture = InputTexture.init(inMTLTexture: resInput, inExpectDim: Dim.init(inDim: expect)) let inputTexture = InputTexture.init(inMTLTexture: resInput, inExpectDim: Dim.init(inDim: dim))
program.scope.setInput(input: inputTexture) program.scope.setInput(input: inputTexture)
//(ops.count - except) //(ops.count - except)
for i in 0..<ops.count { for i in 0..<ops.count {
...@@ -104,9 +104,9 @@ public class Executor<P: PrecisionType> { ...@@ -104,9 +104,9 @@ public class Executor<P: PrecisionType> {
} }
} }
var outputTextures: [Variant]? var outputTextures: [String : [Variant]]?
if except > 0 { if except > 0 {
outputTextures = ops[ops.count - except].inputs() outputTextures = ops[ops.count - except].inputVariant()
} }
buffer.addCompletedHandler { (commandbuffer) in buffer.addCompletedHandler { (commandbuffer) in
...@@ -130,13 +130,11 @@ public class Executor<P: PrecisionType> { ...@@ -130,13 +130,11 @@ public class Executor<P: PrecisionType> {
} }
// return // return
self.ops[testTo].delogOutput() // self.ops[testTo].delogOutput()
// self.ops[91].delogOutput() // self.ops[91].delogOutput()
// self.ops[92].delogOutput() // self.ops[92].delogOutput()
// self.ops[93].delogOutput() // self.ops[93].delogOutput()
return;
let afterDate = Date.init() let afterDate = Date.init()
var resultHolder: ResultHolder<P> var resultHolder: ResultHolder<P>
......
...@@ -25,7 +25,7 @@ protocol Runable { ...@@ -25,7 +25,7 @@ protocol Runable {
func run(device: MTLDevice, buffer: MTLCommandBuffer) throws func run(device: MTLDevice, buffer: MTLCommandBuffer) throws
func runImpl(device: MTLDevice,buffer: MTLCommandBuffer) throws func runImpl(device: MTLDevice,buffer: MTLCommandBuffer) throws
func delogOutput() func delogOutput()
func inputs() -> [Variant] func inputVariant() -> [String : [Variant]]
} }
extension Runable where Self: OperatorProtocol{ extension Runable where Self: OperatorProtocol{
...@@ -35,7 +35,10 @@ extension Runable where Self: OperatorProtocol{ ...@@ -35,7 +35,10 @@ extension Runable where Self: OperatorProtocol{
} catch let error { } catch let error {
throw error throw error
} }
// print(type + ": " + para.outputDesc()) }
func inputVariant() -> [String : [Variant]] {
fatalError(" op \(type) need implement inputVariant")
} }
func delogOutput() { func delogOutput() {
...@@ -98,6 +101,7 @@ class Operator <KernelType: Computable , ParameterType>: OperatorProtocol where ...@@ -98,6 +101,7 @@ class Operator <KernelType: Computable , ParameterType>: OperatorProtocol where
let scope: Scope let scope: Scope
var kernel: KerType var kernel: KerType
required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws {
print("create op: \(opDesc.type)")
type = opDesc.type type = opDesc.type
scope = inScope scope = inScope
inputs = opDesc.inputs inputs = opDesc.inputs
......
...@@ -43,15 +43,11 @@ class BatchNormParam<P: PrecisionType>: OpParam { ...@@ -43,15 +43,11 @@ class BatchNormParam<P: PrecisionType>: OpParam {
} }
class BatchNormOp<P: PrecisionType>: Operator<BatchNormKernel<P>, BatchNormParam<P>>, Runable, Creator, InferShaperable{ class BatchNormOp<P: PrecisionType>: Operator<BatchNormKernel<P>, BatchNormParam<P>>, Runable, Creator, InferShaperable{
typealias OpType = BatchNormOp<P>
func inputs() -> [Variant] {
return [para.input, para.inputBias, para.inputMean, para.inputScale, para.inputVariance]
}
func inferShape() { func inferShape() {
para.output.dim = para.input.dim para.output.dim = para.input.dim
} }
typealias OpType = BatchNormOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
......
///* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. ///* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
// //
// Licensed under the Apache License, Version 2.0 (the "License"); // Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License. // you may not use this file except in compliance with the License.
// You may obtain a copy of the License at // You may obtain a copy of the License at
// //
// http://www.apache.org/licenses/LICENSE-2.0 // http://www.apache.org/licenses/LICENSE-2.0
// //
// Unless required by applicable law or agreed to in writing, software // Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS, // distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and // See the License for the specific language governing permissions and
// limitations under the License. */ // limitations under the License. */
import Foundation
class BoxcoderParam<P: PrecisionType>: OpParam {
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
priorBox = try BoxcoderParam.getFirstTensor(key: "PriorBox", map: opDesc.inputs, from: inScope)
priorBoxVar = try BoxcoderParam.getFirstTensor(key: "PriorBoxVar", map: opDesc.inputs, from: inScope)
targetBox = try BoxcoderParam.getFirstTensor(key: "TargetBox", map: opDesc.inputs, from: inScope)
output = try BoxcoderParam.getFirstTensor(key: "OutputBox", map: opDesc.outputs, from: inScope)
codeType = try BoxcoderParam.getAttr(key: "code_type", attrs: opDesc.attrs)
boxNormalized = try BoxcoderParam.getAttr(key: "box_normalized", attrs: opDesc.attrs)
} catch let error {
throw error
}
assert(priorBox.transpose == [0, 1, 2, 3])
assert(priorBoxVar.transpose == [0, 1, 2, 3])
assert(targetBox.transpose == [0, 1, 2, 3])
assert(codeType == "decode_center_size") // encode_center_size is not implemented
assert((targetBox.tensorDim.cout() == 3) && (targetBox.tensorDim[0] == 1)) // N must be 1 (only handle batch size = 1)
}
let priorBox: Texture<P>
let priorBoxVar: Texture<P>
let targetBox: Texture<P>
var output: Texture<P>
let codeType: String
let boxNormalized: Bool
}
class BoxcoderOp<P: PrecisionType>: Operator<BoxcoderKernel<P>, BoxcoderParam<P>>, Runable, Creator, InferShaperable{
import Foundation typealias OpType = BoxcoderOp<P>
func inferShape() {
// para.output.dim = para.input.dim
}
class BoxcoderParam<P: PrecisionType>: OpParam { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
typealias ParamPrecisionType = P do {
required init(opDesc: OpDesc, inScope: Scope) throws { try kernel.compute(commandBuffer: buffer, param: para)
do { } catch let error {
priorBox = try BoxcoderParam.getFirstTensor(key: "PriorBox", map: opDesc.inputs, from: inScope) throw error
priorBoxVar = try BoxcoderParam.getFirstTensor(key: "PriorBoxVar", map: opDesc.inputs, from: inScope)
targetBox = try BoxcoderParam.getFirstTensor(key: "TargetBox", map: opDesc.inputs, from: inScope)
output = try BoxcoderParam.getFirstTensor(key: "OutputBox", map: opDesc.outputs, from: inScope)
codeType = try BoxcoderParam.getAttr(key: "code_type", attrs: opDesc.attrs)
boxNormalized = try BoxcoderParam.getAttr(key: "box_normalized", attrs: opDesc.attrs)
} catch let error {
throw error
}
assert(priorBox.transpose == [0, 1, 2, 3])
assert(priorBoxVar.transpose == [0, 1, 2, 3])
assert(targetBox.transpose == [0, 1, 2, 3])
assert(codeType == "decode_center_size") // encode_center_size is not implemented
assert((targetBox.tensorDim.cout() == 3) && (targetBox.tensorDim[0] == 1)) // N must be 1 (only handle batch size = 1)
} }
let priorBox: Texture<P>
let priorBoxVar: Texture<P>
let targetBox: Texture<P>
var output: Texture<P>
let codeType: String
let boxNormalized: Bool
} }
class BoxcoderOp<P: PrecisionType>: Operator<BoxcoderKernel<P>, BoxcoderParam<P>>, Runable, Creator, InferShaperable{ func delogOutput() {
func inputs() -> [Variant] { print(" \(type) output: ")
return [para.priorBox, para.priorBoxVar, para.targetBox] 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())
func inferShape() { let priorBoxVarOriginDim = para.priorBoxVar.originDim
// para.output.dim = para.input.dim 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())
typealias OpType = BoxcoderOp<P> let targetBoxOriginDim = para.targetBox.originDim
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { let targetBoxArray = para.targetBox.metalTexture.realNHWC(dim: (n: targetBoxOriginDim[0], h: targetBoxOriginDim[1], w: targetBoxOriginDim[2], c: targetBoxOriginDim[3]))
do { print(" target box ")
try kernel.compute(commandBuffer: buffer, param: para) print(targetBoxArray.strideArray())
} catch let error {
throw error
}
}
func delogOutput() { let originDim = para.output.originDim
// let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in
// return o let outputArray = para.output.metalTexture.realNHWC(dim: (n: originDim[0], h: originDim[1], w: originDim[2], c: originDim[3]))
// } print(outputArray.strideArray())
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 ")
}
} }
}
...@@ -40,16 +40,13 @@ class ConcatParam<P: PrecisionType>: OpParam { ...@@ -40,16 +40,13 @@ class ConcatParam<P: PrecisionType>: OpParam {
class ConcatOp<P: PrecisionType>: Operator<ConcatKernel<P>, ConcatParam<P>>, Runable, Creator, InferShaperable{ class ConcatOp<P: PrecisionType>: Operator<ConcatKernel<P>, ConcatParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] { typealias OpType = ConcatOp<P>
return para.input
}
func inferShape() { func inferShape() {
// let dim = para.input.reduce([0, 0]) {[$0[0] + $1.dim[0], $1.dim[1]]} // let dim = para.input.reduce([0, 0]) {[$0[0] + $1.dim[0], $1.dim[1]]}
// para.output.dim = Dim.init(inDim: dim) // para.output.dim = Dim.init(inDim: dim)
} }
typealias OpType = ConcatOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
...@@ -59,34 +56,10 @@ class ConcatOp<P: PrecisionType>: Operator<ConcatKernel<P>, ConcatParam<P>>, Run ...@@ -59,34 +56,10 @@ class ConcatOp<P: PrecisionType>: Operator<ConcatKernel<P>, ConcatParam<P>>, Run
} }
func delogOutput() { func delogOutput() {
print(" \(type) output: ")
let originDim = para.output.originDim let originDim = para.output.originDim
let outputArray = para.output.metalTexture.realNHWC(dim: (n: originDim[0], h: originDim[1], w: originDim[2], c: originDim[3])) 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())
// let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in
// return o
// }
//
//// print(outputArray.strideArray())
//
// writeToLibrary(fileName: "concat_out", array: outputArray)
//
// let device: MTLDevice = MTLCreateSystemDefaultDevice()!
// let tensorArray: [P] = device.texture2tensor(texture: para.output.metalTexture, dim: [1917, 4])
// print(tensorArray.strideArray())
// print(para.output.metalTexture)
// writeToLibrary(fileName: "concat_out", array: outputArray)
// print(" write done ")
// print(outputArray.strideArray())
} }
} }
......
...@@ -61,10 +61,6 @@ class ConvAddBatchNormReluParam<P: PrecisionType>: OpParam { ...@@ -61,10 +61,6 @@ class ConvAddBatchNormReluParam<P: PrecisionType>: OpParam {
class ConvAddBatchNormReluOp<P: PrecisionType>: Operator<ConvAddBatchNormReluKernel<P>, ConvAddBatchNormReluParam<P>>, Runable, Creator, InferShaperable, Fusion{ class ConvAddBatchNormReluOp<P: PrecisionType>: Operator<ConvAddBatchNormReluKernel<P>, ConvAddBatchNormReluParam<P>>, Runable, Creator, InferShaperable, Fusion{
func inputs() -> [Variant] {
return [para.variance, para.bias, para.mean, para.scale, para.y, para.filter, para.input]
}
typealias OpType = ConvAddBatchNormReluOp<P> typealias OpType = ConvAddBatchNormReluOp<P>
func inferShape() { func inferShape() {
...@@ -115,7 +111,8 @@ class ConvAddBatchNormReluOp<P: PrecisionType>: Operator<ConvAddBatchNormReluKer ...@@ -115,7 +111,8 @@ class ConvAddBatchNormReluOp<P: PrecisionType>: Operator<ConvAddBatchNormReluKer
} }
func delogOutput() { func delogOutput() {
print(" conv add batchnorm relu output ")
print(para.output.metalTexture.toTensor(dim: (n: para.output.originDim[0], c: para.output.originDim[1], h: para.output.originDim[2], w: para.output.originDim[3])).strideArray())
// let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false) // let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false)
// para.filter.logDataPointer(header: "filter data pointer: ") // para.filter.logDataPointer(header: "filter data pointer: ")
// print("filter: \(para.filter)") // print("filter: \(para.filter)")
......
...@@ -43,30 +43,9 @@ class ConvAddParam<P: PrecisionType>: OpParam { ...@@ -43,30 +43,9 @@ class ConvAddParam<P: PrecisionType>: OpParam {
} }
class ConvAddOp<P: PrecisionType>: Operator<ConvAddKernel<P>, ConvAddParam<P>>, Runable, Creator, InferShaperable, Fusion{ class ConvAddOp<P: PrecisionType>: Operator<ConvAddKernel<P>, ConvAddParam<P>>, Runable, Creator, InferShaperable, Fusion{
typealias OpType = ConvAddOp<P>
func delogOutput() {
print(para.output.metalTexture.toTensor(dim: (n: para.output.tensorDim[0], c: para.output.tensorDim[1], h: para.output.tensorDim[2], w: para.output.tensorDim[3])).strideArray())
// print(" conv add: ")
// print(para.input.metalTexture)
// print(" filter array: ")
// let filterArray: [P] = para.filter.buffer.array()
// print(filterArray)
// let input = para.input.metalTexture.floatArray { (p: P) -> P in
// return p
// }
// print(input)
// let output = para.output.metalTexture.floatArray { (p: P) -> P in
// return p
// }
// print(para.output.metalTexture)
// print(output)
}
static func fusionNode() -> Node { static func fusionNode() -> Node {
...@@ -80,16 +59,10 @@ class ConvAddOp<P: PrecisionType>: Operator<ConvAddKernel<P>, ConvAddParam<P>>, ...@@ -80,16 +59,10 @@ class ConvAddOp<P: PrecisionType>: Operator<ConvAddKernel<P>, ConvAddParam<P>>,
return [:] return [:]
} }
func inputs() -> [Variant] {
return [para.input, para.y, para.filter]
}
static func fusionType() -> String { static func fusionType() -> String {
return gConvAddType return gConvAddType
} }
typealias OpType = ConvAddOp<P>
func inferShape() { func inferShape() {
...@@ -122,4 +95,7 @@ class ConvAddOp<P: PrecisionType>: Operator<ConvAddKernel<P>, ConvAddParam<P>>, ...@@ -122,4 +95,7 @@ class ConvAddOp<P: PrecisionType>: Operator<ConvAddKernel<P>, ConvAddParam<P>>,
} }
} }
func delogOutput() {
print(para.output.metalTexture.toTensor(dim: (n: para.output.tensorDim[0], c: para.output.tensorDim[1], h: para.output.tensorDim[2], w: para.output.tensorDim[3])).strideArray())
}
} }
...@@ -109,74 +109,8 @@ class ConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNReluPa ...@@ -109,74 +109,8 @@ class ConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNReluPa
} }
func delogOutput() { func delogOutput() {
print(" \(type) output: ")
// let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false) print(para.output.metalTexture.toTensor(dim: (n: para.output.originDim[0], c: para.output.originDim[1], h: para.output.originDim[2], w: para.output.originDim[3])).strideArray())
// 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)
// print("input: ")
// print(para.input.metalTexture)
//
// let input = para.input.metalTexture.floatArray { (p: P) -> P in
// return p
// }
// for i in 0..<input.count {
// print(" index \(i) : \(input[i])")
// }
// print(input)
// writeToLibrary(fileName: "input35", array: input)
// print(input)
print(para.newBiase?.length)
print(para.newScale?.length)
// let newScale = para.newScale?.contents().bindMemory(to: P.self, capacity: para.newScale!.length)
// let newBiase = para.newBiase?.contents().bindMemory(to: P.self, capacity: para.newBiase!.length)
//
// let filterArray: [Float32] = para.filter.buffer.array();
//// writeToLibrary(fileName: "filter35", array: filterArray)
//
// print(filterArray)
//
// print("new scale: ")
// for i in 0..<(para.newScale!.length / MemoryLayout<P>.size) {
// print("index: \(i) \(newScale![i]) ")
// }
//
// print("new biase: ")
// for i in 0..<(para.newBiase!.length / MemoryLayout<P>.size) {
// print("index: \(i) \(newBiase![i]) ")
// }
// print(para.output.metalTexture)
//
//
//
// let output = para.output.metalTexture.floatArray { (p: P) -> P in
// return p
// }
// print(output)
//
//
// writeToLibrary(fileName: "batch_norm_34.tmp_2", array: output)
// print(" write done")
//
print(para.output.metalTexture.toTensor(dim: (n: para.output.tensorDim[0], c: para.output.tensorDim[1], h: para.output.tensorDim[2], w: para.output.tensorDim[3])).strideArray())
// let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: true)
} }
} }
...@@ -41,19 +41,8 @@ class ConvParam<P: PrecisionType>: OpParam { ...@@ -41,19 +41,8 @@ class ConvParam<P: PrecisionType>: OpParam {
} }
class ConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable, Creator, InferShaperable { class ConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable, Creator, InferShaperable {
typealias OpType = ConvOp<P>
func inputs() -> [Variant] {
return [para.input, para.filter]
}
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() { func inferShape() {
let inDims = para.input.dim let inDims = para.input.dim
let filterDim = para.filter.dim let filterDim = para.filter.dim
...@@ -76,7 +65,6 @@ class ConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable, ...@@ -76,7 +65,6 @@ class ConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable,
para.output.dim = Dim.init(inDim: outDim) para.output.dim = Dim.init(inDim: outDim)
} }
typealias OpType = ConvOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
...@@ -87,7 +75,7 @@ class ConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable, ...@@ -87,7 +75,7 @@ class ConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable,
func delogOutput() { func delogOutput() {
print("conv output : ") print("conv output : ")
print(para.output.metalTexture) print(para.output.metalTexture.toTensor(dim: (n: para.output.originDim[0], c: para.output.originDim[1], h: para.output.originDim[2], w: para.output.originDim[3])).strideArray())
// let _: Float16? = para.output.metalTexture.logDesc() // let _: Float16? = para.output.metalTexture.logDesc()
} }
} }
...@@ -28,15 +28,12 @@ class ConvTransposeParam<P: PrecisionType>: ConvParam<P> { ...@@ -28,15 +28,12 @@ class ConvTransposeParam<P: PrecisionType>: ConvParam<P> {
class ConvTransposeOp<P: PrecisionType>: Operator<ConvTransposeKernel<P>, ConvTransposeParam<P>>, Runable, Creator, InferShaperable{ class ConvTransposeOp<P: PrecisionType>: Operator<ConvTransposeKernel<P>, ConvTransposeParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] { typealias OpType = ConvTransposeOp<P>
return [para.input, para.filter]
}
func inferShape() { func inferShape() {
// para.output.dim = para.input.dim // para.output.dim = para.input.dim
} }
typealias OpType = ConvTransposeOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
...@@ -44,6 +41,7 @@ class ConvTransposeOp<P: PrecisionType>: Operator<ConvTransposeKernel<P>, ConvTr ...@@ -44,6 +41,7 @@ class ConvTransposeOp<P: PrecisionType>: Operator<ConvTransposeKernel<P>, ConvTr
throw error throw error
} }
} }
func delogOutput() { func delogOutput() {
print("conv transpose delog") print("conv transpose delog")
let _: P? = para.input.metalTexture.logDesc(header: "conv transpose input: ", stridable: true) let _: P? = para.input.metalTexture.logDesc(header: "conv transpose input: ", stridable: true)
......
...@@ -15,11 +15,9 @@ ...@@ -15,11 +15,9 @@
import Foundation import Foundation
class DepthConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable, Creator, InferShaperable { class DepthConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable, Creator, InferShaperable {
func inputs() -> [Variant] { typealias OpType = DepthConvOp<P>
return [para.input, para.filter]
}
required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws {
do { do {
try super.init(device: device, opDesc: opDesc, inScope: inScope) try super.init(device: device, opDesc: opDesc, inScope: inScope)
...@@ -50,8 +48,6 @@ class DepthConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runa ...@@ -50,8 +48,6 @@ class DepthConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runa
para.output.dim = Dim.init(inDim: outDim) para.output.dim = Dim.init(inDim: outDim)
} }
typealias OpType = DepthConvOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
...@@ -61,8 +57,7 @@ class DepthConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runa ...@@ -61,8 +57,7 @@ class DepthConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runa
} }
func delogOutput() { func delogOutput() {
print("conv output : ") print(" \(type) output: ")
print(para.output.metalTexture) print(para.output.metalTexture.toTensor(dim: (n: para.output.originDim[0], c: para.output.originDim[1], h: para.output.originDim[2], w: para.output.originDim[3])).strideArray())
// let _: Float16? = para.output.metalTexture.logDesc()
} }
} }
...@@ -17,10 +17,6 @@ import Foundation ...@@ -17,10 +17,6 @@ import Foundation
class DwConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNReluParam<P>>, Runable, Creator, InferShaperable, Fusion{ class DwConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNReluParam<P>>, Runable, Creator, InferShaperable, Fusion{
typealias OpType = ConvBNReluOp<P> typealias OpType = ConvBNReluOp<P>
func inputs() -> [Variant] {
return [para.input, para.bias, para.mean, para.filter, para.variance, para.scale]
}
func inferShape() { func inferShape() {
let inDims = para.input.dim let inDims = para.input.dim
let filterDim = para.filter.dim let filterDim = para.filter.dim
...@@ -68,26 +64,7 @@ class DwConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNRelu ...@@ -68,26 +64,7 @@ class DwConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNRelu
} }
func delogOutput() { func delogOutput() {
print(" \(type) output: ")
// let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false) print(para.output.metalTexture.toTensor(dim: (n: para.output.originDim[0], c: para.output.originDim[1], h: para.output.originDim[2], w: para.output.originDim[3])).strideArray())
// 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: "batch_norm_19.tmp_2", array: output)
// print(" write done")
// let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false)
} }
} }
...@@ -34,16 +34,12 @@ class ElementwiseAddParam<P: PrecisionType>: OpParam { ...@@ -34,16 +34,12 @@ class ElementwiseAddParam<P: PrecisionType>: OpParam {
} }
class ElementwiseAddOp<P: PrecisionType>: Operator<ElementwiseAddKernel<P>, ElementwiseAddParam<P>>, Runable, Creator, InferShaperable{ class ElementwiseAddOp<P: PrecisionType>: Operator<ElementwiseAddKernel<P>, ElementwiseAddParam<P>>, Runable, Creator, InferShaperable{
typealias OpType = ElementwiseAddOp<P>
func inputs() -> [Variant] {
return [para.input, para.inputY]
}
func inferShape() { func inferShape() {
para.output.dim = para.input.dim para.output.dim = para.input.dim
} }
typealias OpType = ElementwiseAddOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
} }
} }
......
...@@ -35,11 +35,7 @@ class FeedParam<P: PrecisionType>: OpParam{ ...@@ -35,11 +35,7 @@ class FeedParam<P: PrecisionType>: OpParam{
class FeedOp<P: PrecisionType>: Operator<Texture2DTo2DArrayKernel<P>, FeedParam<P>>, Runable, Creator, InferShaperable { class FeedOp<P: PrecisionType>: Operator<Texture2DTo2DArrayKernel<P>, FeedParam<P>>, Runable, Creator, InferShaperable {
typealias OpType = FeedOp<P> typealias OpType = FeedOp<P>
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() { func inferShape() {
// print("feed input: \(para.input.expectDim)") // print("feed input: \(para.input.expectDim)")
print("feed output: \(para.output.dim)") print("feed output: \(para.output.dim)")
......
...@@ -37,20 +37,18 @@ class FetchKernel<P: PrecisionType>: Kernel, Computable { ...@@ -37,20 +37,18 @@ class FetchKernel<P: PrecisionType>: Kernel, Computable {
} }
required init(device: MTLDevice, param: FetchParam<P>) { required init(device: MTLDevice, param: FetchParam<P>) {
super.init(device: device, inFunctionName: "texture2d_to_2d_array") super.init(device: device, inFunctionName: "place_holder")
} }
} }
class FetchOp<P: PrecisionType>: Operator< FetchKernel<P>, FetchParam<P>>, Runable, Creator, InferShaperable{ class FetchOp<P: PrecisionType>: Operator< FetchKernel<P>, FetchParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input]
}
typealias OpType = FetchOp<P>
func inferShape() { func inferShape() {
print(para.input.dim) print(para.input.dim)
} }
typealias OpType = FetchOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
scope.setOutput(output: para.output) scope.setOutput(output: para.output)
} }
......
...@@ -43,7 +43,6 @@ class ConvTransposeKernel<P: PrecisionType>: Kernel, Computable{ ...@@ -43,7 +43,6 @@ class ConvTransposeKernel<P: PrecisionType>: Kernel, Computable{
let dilationY = UInt16(param.dilations[1]) let dilationY = UInt16(param.dilations[1])
metalParam = MetalConvTransposeParam.init(kernelW: kernelWidth, kernelH: kernelHeight, strideX: strideX, strideY: strideY, paddingX: paddingX, paddingY: paddingY, dilationX: dilationX, dilationY: dilationY) metalParam = MetalConvTransposeParam.init(kernelW: kernelWidth, kernelH: kernelHeight, strideX: strideX, strideY: strideY, paddingX: paddingX, paddingY: paddingY, dilationX: dilationX, dilationY: dilationY)
} }
func compute(commandBuffer: MTLCommandBuffer, param: ConvTransposeParam<P>) throws { func compute(commandBuffer: MTLCommandBuffer, param: ConvTransposeParam<P>) throws {
......
...@@ -15,17 +15,11 @@ ...@@ -15,17 +15,11 @@
import Foundation import Foundation
class MulticlassNMSKernel<P: PrecisionType>: Kernel, Computable{ class MulticlassNMSKernel<P: PrecisionType>: Kernel, Computable{
func compute(commandBuffer: MTLCommandBuffer, param: MulticlassNMSParam<P>) throws {
guard let encoder = commandBuffer.makeComputeCommandEncoder() else { required init(device: MTLDevice, param: MulticlassNMSParam<P>) {
throw PaddleMobileError.predictError(message: " encode is nil") super.init(device: device, inFunctionName: "place_holder")
} }
// encoder.setTexture(param.input.metalTexture, index: 0)
// encoder.setTexture(param.output.metalTexture, index: 1) func compute(commandBuffer: MTLCommandBuffer, param: MulticlassNMSParam<P>) throws {
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) }
encoder.endEncoding()
}
required init(device: MTLDevice, param: MulticlassNMSParam<P>) {
super.init(device: device, inFunctionName: "prior_box")
}
} }
...@@ -15,46 +15,48 @@ ...@@ -15,46 +15,48 @@
import Foundation import Foundation
struct PoolMetalParam { struct PoolMetalParam {
let ksizeX: Int32 let ksizeX: Int32
let ksizeY: Int32 let ksizeY: Int32
let strideX: Int32 let strideX: Int32
let strideY: Int32 let strideY: Int32
let paddingX: Int32 let paddingX: Int32
let paddingY: Int32 let paddingY: Int32
let poolType: Int32 let poolType: Int32
} }
class PoolKernel<P: PrecisionType>: Kernel, Computable{ class PoolKernel<P: PrecisionType>: Kernel, Computable{
func compute(commandBuffer: MTLCommandBuffer, param: PoolParam<P>) throws {
guard let encoder = commandBuffer.makeComputeCommandEncoder() else { required init(device: MTLDevice, param: PoolParam<P>) {
throw PaddleMobileError.predictError(message: " encoder is nil") super.init(device: device, inFunctionName: "pool")
} param.output.initTexture(device: device)
encoder.setTexture(param.input.metalTexture, index: 0) }
encoder.setTexture(param.output.metalTexture, index: 1)
var poolType: Int32 func compute(commandBuffer: MTLCommandBuffer, param: PoolParam<P>) throws {
switch param.poolType { guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
case "max": throw PaddleMobileError.predictError(message: " encoder is nil")
poolType = 0
case "avg":
poolType = 1
default:
throw PaddleMobileError.predictError(message: " unknown pooltype " + param.poolType)
}
var pmp = PoolMetalParam.init(
ksizeX: param.ksize[0],
ksizeY: param.ksize[1],
strideX: param.stride[0],
strideY: param.stride[1],
paddingX: param.padding[0],
paddingY: param.padding[1],
poolType: poolType
)
encoder.setBytes(&pmp, length: MemoryLayout<PoolMetalParam>.size, index: 0)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
} }
encoder.setTexture(param.input.metalTexture, index: 0)
required init(device: MTLDevice, param: PoolParam<P>) { encoder.setTexture(param.output.metalTexture, index: 1)
super.init(device: device, inFunctionName: "pool") var poolType: Int32
switch param.poolType {
case "max":
poolType = 0
case "avg":
poolType = 1
default:
throw PaddleMobileError.predictError(message: " unknown pooltype " + param.poolType)
} }
var pmp = PoolMetalParam.init(
ksizeX: param.ksize[0],
ksizeY: param.ksize[1],
strideX: param.stride[0],
strideY: param.stride[1],
paddingX: param.padding[0],
paddingY: param.padding[1],
poolType: poolType
)
encoder.setBytes(&pmp, length: MemoryLayout<PoolMetalParam>.size, index: 0)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
}
} }
// /* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
// common.metal
// paddle-mobile Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// Created by liuRuiLong on 2018/8/26. You may obtain a copy of the License at
// Copyright © 2018年 orange. All rights reserved.
// 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> #include <metal_stdlib>
using namespace metal; using namespace metal;
......
...@@ -16,6 +16,12 @@ ...@@ -16,6 +16,12 @@
#include "Common.metal" #include "Common.metal"
using namespace metal; using namespace metal;
// 占位函数, 啥也没干
kernel void place_holder(texture2d<half, access::read> inTexture [[texture(0)]],
texture2d_array<half, access::write> outTexture [[texture(1)]],
uint3 gid [[thread_position_in_grid]]) {
}
struct OutputDim { struct OutputDim {
ushort width; ushort width;
ushort height; ushort height;
...@@ -161,9 +167,6 @@ kernel void pool_half(texture2d_array<half, access::read> inTexture [[texture(0) ...@@ -161,9 +167,6 @@ kernel void pool_half(texture2d_array<half, access::read> inTexture [[texture(0)
outTexture.write(r, gid.xy, gid.z); outTexture.write(r, gid.xy, gid.z);
} }
struct TransposeParam { struct TransposeParam {
int iC; int iC;
int oC; int oC;
......
...@@ -15,39 +15,39 @@ ...@@ -15,39 +15,39 @@
import Foundation import Foundation
class MulticlassNMSParam<P: PrecisionType>: OpParam { class MulticlassNMSParam<P: PrecisionType>: OpParam {
typealias ParamPrecisionType = P typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws { required init(opDesc: OpDesc, inScope: Scope) throws {
do { do {
scores = try MulticlassNMSParam.getFirstTensor(key: "Scores", map: opDesc.inputs, from: inScope) scores = try MulticlassNMSParam.getFirstTensor(key: "Scores", map: opDesc.inputs, from: inScope)
bboxes = try MulticlassNMSParam.getFirstTensor(key: "BBoxes", map: opDesc.inputs, from: inScope) bboxes = try MulticlassNMSParam.getFirstTensor(key: "BBoxes", map: opDesc.inputs, from: inScope)
output = try MulticlassNMSParam.outputOut(outputs: opDesc.outputs, from: inScope) output = try MulticlassNMSParam.outputOut(outputs: opDesc.outputs, from: inScope)
} catch let error { } catch let error {
throw error throw error
}
} }
let scores: Texture<P> }
let bboxes: Texture<P> let scores: Texture<P>
var output: Texture<P> let bboxes: Texture<P>
var output: Texture<P>
} }
class MulticlassNMSOp<P: PrecisionType>: Operator<MulticlassNMSKernel<P>, MulticlassNMSParam<P>>, Runable, Creator, InferShaperable{ class MulticlassNMSOp<P: PrecisionType>: Operator<MulticlassNMSKernel<P>, MulticlassNMSParam<P>>, Runable, Creator, InferShaperable{
func inputVariant() -> [String : [Variant]] {
return ["Scores" : [para.scores], "BBoxes" : [para.bboxes]]
}
func inputs() -> [Variant] { func inferShape() {
return [para.scores,para.bboxes] // para.output.dim = para.input.dim
} }
func inferShape() { typealias OpType = MulticlassNMSOp<P>
// para.output.dim = para.input.dim func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
} do {
try kernel.compute(commandBuffer: buffer, param: para)
typealias OpType = MulticlassNMSOp<P> } catch let error {
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { throw error
do {
try kernel.compute(commandBuffer: buffer, param: para)
} catch let error {
throw error
}
} }
}
} }
......
...@@ -43,15 +43,12 @@ class PoolParam<P: PrecisionType>: OpParam { ...@@ -43,15 +43,12 @@ class PoolParam<P: PrecisionType>: OpParam {
class PoolOp<P: PrecisionType>: Operator<PoolKernel<P>, PoolParam<P>>, Runable, Creator, InferShaperable{ class PoolOp<P: PrecisionType>: Operator<PoolKernel<P>, PoolParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] { typealias OpType = PoolOp<P>
return [para.input]
}
func inferShape() { func inferShape() {
// para.output.dim = para.input.dim // para.output.dim = para.input.dim
} }
typealias OpType = PoolOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
......
...@@ -35,15 +35,12 @@ class PreluParam<P: PrecisionType>: OpParam { ...@@ -35,15 +35,12 @@ class PreluParam<P: PrecisionType>: OpParam {
class PreluOp<P: PrecisionType>: Operator<PreluKernel<P>, PreluParam<P>>, Runable, Creator, InferShaperable{ class PreluOp<P: PrecisionType>: Operator<PreluKernel<P>, PreluParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] { typealias OpType = PreluOp<P>
return [para.alpha, para.input]
}
func inferShape() { func inferShape() {
// para.output.dim = para.input.dim // para.output.dim = para.input.dim
} }
typealias OpType = PreluOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
...@@ -51,6 +48,7 @@ class PreluOp<P: PrecisionType>: Operator<PreluKernel<P>, PreluParam<P>>, Runabl ...@@ -51,6 +48,7 @@ class PreluOp<P: PrecisionType>: Operator<PreluKernel<P>, PreluParam<P>>, Runabl
throw error throw error
} }
} }
func delogOutput() { func delogOutput() {
print("softmax delog") print("softmax delog")
let _: P? = para.input.metalTexture.logDesc(header: "softmax input: ", stridable: false) let _: P? = para.input.metalTexture.logDesc(header: "softmax input: ", stridable: false)
......
...@@ -55,15 +55,11 @@ class PriorBoxParam<P: PrecisionType>: OpParam { ...@@ -55,15 +55,11 @@ class PriorBoxParam<P: PrecisionType>: OpParam {
class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P>>, Runable, Creator, InferShaperable{ class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] { typealias OpType = PriorBoxOp<P>
return [para.input, para.inputImage]
}
func inferShape() { func inferShape() {
} }
typealias OpType = PriorBoxOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
...@@ -73,8 +69,8 @@ class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P> ...@@ -73,8 +69,8 @@ class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P>
} }
func delogOutput() { func delogOutput() {
print("pribox: ") print(" \(type) output: ")
print("output: ")
// output // output
let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in
return o return o
...@@ -93,9 +89,6 @@ class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P> ...@@ -93,9 +89,6 @@ class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P>
// writeToLibrary(fileName: "variance_out", array: outputVarianceArray) // writeToLibrary(fileName: "variance_out", array: outputVarianceArray)
print("pribox write done ")
} }
} }
......
...@@ -30,15 +30,12 @@ class ReluParam<P: PrecisionType>: OpParam { ...@@ -30,15 +30,12 @@ class ReluParam<P: PrecisionType>: OpParam {
class ReluOp<P: PrecisionType>: Operator<ReluKernel<P>, ReluParam<P>>, Runable, Creator, InferShaperable{ class ReluOp<P: PrecisionType>: Operator<ReluKernel<P>, ReluParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] { typealias OpType = ReluOp<P>
return [para.input]
}
func inferShape() { func inferShape() {
para.output.dim = para.input.dim para.output.dim = para.input.dim
} }
typealias OpType = ReluOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
......
...@@ -57,15 +57,12 @@ class ReshapeParam<P: PrecisionType>: OpParam { ...@@ -57,15 +57,12 @@ class ReshapeParam<P: PrecisionType>: OpParam {
class ReshapeOp<P: PrecisionType>: Operator<ReshapeKernel<P>, ReshapeParam<P>>, Runable, Creator, InferShaperable{ class ReshapeOp<P: PrecisionType>: Operator<ReshapeKernel<P>, ReshapeParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] { typealias OpType = ReshapeOp<P>
return [para.input]
}
func inferShape() { func inferShape() {
// para.output.dim = para.input.dim // para.output.dim = para.input.dim
} }
typealias OpType = ReshapeOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
......
...@@ -36,16 +36,12 @@ class SoftmaxParam<P: PrecisionType>: OpParam { ...@@ -36,16 +36,12 @@ class SoftmaxParam<P: PrecisionType>: OpParam {
} }
class SoftmaxOp<P: PrecisionType>: Operator<SoftmaxKernel<P>, SoftmaxParam<P>>, Runable, Creator, InferShaperable{ class SoftmaxOp<P: PrecisionType>: Operator<SoftmaxKernel<P>, SoftmaxParam<P>>, Runable, Creator, InferShaperable{
typealias OpType = SoftmaxOp<P>
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() { func inferShape() {
// para.output.dim = para.input.dim // para.output.dim = para.input.dim
} }
typealias OpType = SoftmaxOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
...@@ -53,6 +49,7 @@ class SoftmaxOp<P: PrecisionType>: Operator<SoftmaxKernel<P>, SoftmaxParam<P>>, ...@@ -53,6 +49,7 @@ class SoftmaxOp<P: PrecisionType>: Operator<SoftmaxKernel<P>, SoftmaxParam<P>>,
throw error throw error
} }
} }
func delogOutput() { func delogOutput() {
print("softmax delog") print("softmax delog")
......
...@@ -32,15 +32,12 @@ class TransposeParam<P: PrecisionType>: OpParam { ...@@ -32,15 +32,12 @@ class TransposeParam<P: PrecisionType>: OpParam {
class TransposeOp<P: PrecisionType>: Operator<TransposeKernel<P>, TransposeParam<P>>, Runable, Creator, InferShaperable{ class TransposeOp<P: PrecisionType>: Operator<TransposeKernel<P>, TransposeParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] { typealias OpType = TransposeOp<P>
return [para.input]
}
func inferShape() { func inferShape() {
//para.output.dim = para.input.dim //para.output.dim = para.input.dim
} }
typealias OpType = TransposeOp<P>
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do { do {
try kernel.compute(commandBuffer: buffer, param: para) try kernel.compute(commandBuffer: buffer, param: para)
...@@ -48,32 +45,13 @@ class TransposeOp<P: PrecisionType>: Operator<TransposeKernel<P>, TransposeParam ...@@ -48,32 +45,13 @@ class TransposeOp<P: PrecisionType>: Operator<TransposeKernel<P>, TransposeParam
throw error throw error
} }
} }
func delogOutput() { func delogOutput() {
print(para.output.metalTexture.realNHWC(dim: (n: para.output.originDim[0], h: para.output.originDim[1], w: para.output.originDim[2], c: para.output.originDim[3])).strideArray()) print(" \(type) output: ")
let originDim = para.output.tensorDim
// print(para.input.metalTexture.toTensor(dim: (n: para.input.originDim[0], c: para.input.originDim[1], h: para.input.originDim[2], w: para.input.originDim[3])).strideArray()) let outputArray = para.output.metalTexture.realNHWC(dim: (n: originDim[0], h: originDim[1], w: originDim[2], c: originDim[3]))
// print(outputArray.strideArray())
//
// let originDim = para.output.tensorDim
// let outputArray = para.output.metalTexture.realNHWC(dim: (n: originDim[0], h: originDim[1], w: originDim[2], c: originDim[3]))
// print(outputArray.strideArray())
// let inputArray: [Float32] = para.input.metalTexture.floatArray { (ele: Float32) -> Float32 in
// return ele
// }
//
// print(inputArray.strideArray())
//
// let outputArray: [Float32] = para.output.metalTexture.floatArray { (ele: Float32) -> Float32 in
// return ele
// }
//
// print(outputArray.strideArray())
// writeToLibrary(fileName: "transpose_ouput", array: outputArray)
} }
} }
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册