提交 868e1eee 编写于 作者: L liuruilong

support super resolution

上级 9e6ce432
......@@ -17,10 +17,12 @@
FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8820E11C560081E9F8 /* Assets.xcassets */; };
FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */; };
FC203FB221CBFDBA00B37166 /* test.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FC203FA921CBFDBA00B37166 /* test.jpg */; };
FC203FB321CBFDBA00B37166 /* combined_mobilenet_params in Resources */ = {isa = PBXBuildFile; fileRef = FC203FAD21CBFDBA00B37166 /* combined_mobilenet_params */; };
FC203FB421CBFDBA00B37166 /* combined_mobilenet_model in Resources */ = {isa = PBXBuildFile; fileRef = FC203FAE21CBFDBA00B37166 /* combined_mobilenet_model */; };
FC203FB521CBFDBA00B37166 /* yolo_params in Resources */ = {isa = PBXBuildFile; fileRef = FC203FB021CBFDBA00B37166 /* yolo_params */; };
FC203FB621CBFDBA00B37166 /* yolo_model in Resources */ = {isa = PBXBuildFile; fileRef = FC203FB121CBFDBA00B37166 /* yolo_model */; };
FC704C1921D2375300F98BAB /* super_params in Resources */ = {isa = PBXBuildFile; fileRef = FC704C1721D2375300F98BAB /* super_params */; };
FC704C1A21D2375300F98BAB /* super_model in Resources */ = {isa = PBXBuildFile; fileRef = FC704C1821D2375300F98BAB /* super_model */; };
FC704C2221D237FC00F98BAB /* combined_mobilenet_params in Resources */ = {isa = PBXBuildFile; fileRef = FC704C1D21D237FC00F98BAB /* combined_mobilenet_params */; };
FC704C2321D237FC00F98BAB /* combined_mobilenet_model in Resources */ = {isa = PBXBuildFile; fileRef = FC704C1E21D237FC00F98BAB /* combined_mobilenet_model */; };
FC704C2421D237FC00F98BAB /* yolo_params in Resources */ = {isa = PBXBuildFile; fileRef = FC704C2021D237FC00F98BAB /* yolo_params */; };
FC704C2521D237FC00F98BAB /* yolo_model in Resources */ = {isa = PBXBuildFile; fileRef = FC704C2121D237FC00F98BAB /* yolo_model */; };
FC803BCD214D27930094B8E5 /* FPSCounter.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC803BCB214D27920094B8E5 /* FPSCounter.swift */; };
FC803BCE214D27930094B8E5 /* VideoCapture.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC803BCC214D27920094B8E5 /* VideoCapture.swift */; };
FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC542122EF5400D94F7E /* MetalHelper.swift */; };
......@@ -58,12 +60,14 @@
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>"; };
FC203FA921CBFDBA00B37166 /* test.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = test.jpg; sourceTree = "<group>"; };
FC203FAD21CBFDBA00B37166 /* combined_mobilenet_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = combined_mobilenet_params; sourceTree = "<group>"; };
FC203FAE21CBFDBA00B37166 /* combined_mobilenet_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = combined_mobilenet_model; sourceTree = "<group>"; };
FC203FB021CBFDBA00B37166 /* yolo_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = yolo_params; sourceTree = "<group>"; };
FC203FB121CBFDBA00B37166 /* yolo_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = yolo_model; sourceTree = "<group>"; };
FC27991121343A39000B6BAD /* paddle-mobile-demo-Bridging-Header.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = "paddle-mobile-demo-Bridging-Header.h"; sourceTree = "<group>"; };
FC4FD97B2140EE250073E130 /* libc++.tbd */ = {isa = PBXFileReference; lastKnownFileType = "sourcecode.text-based-dylib-definition"; name = "libc++.tbd"; path = "usr/lib/libc++.tbd"; sourceTree = SDKROOT; };
FC704C1721D2375300F98BAB /* super_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = super_params; sourceTree = "<group>"; };
FC704C1821D2375300F98BAB /* super_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = super_model; sourceTree = "<group>"; };
FC704C1D21D237FC00F98BAB /* combined_mobilenet_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = combined_mobilenet_params; sourceTree = "<group>"; };
FC704C1E21D237FC00F98BAB /* combined_mobilenet_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = combined_mobilenet_model; sourceTree = "<group>"; };
FC704C2021D237FC00F98BAB /* yolo_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = yolo_params; sourceTree = "<group>"; };
FC704C2121D237FC00F98BAB /* yolo_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = yolo_model; sourceTree = "<group>"; };
FC803BCB214D27920094B8E5 /* FPSCounter.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = FPSCounter.swift; sourceTree = "<group>"; };
FC803BCC214D27920094B8E5 /* VideoCapture.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = VideoCapture.swift; sourceTree = "<group>"; };
FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = "<group>"; };
......@@ -155,35 +159,45 @@
FC203FAA21CBFDBA00B37166 /* models */ = {
isa = PBXGroup;
children = (
FC203FAB21CBFDBA00B37166 /* vision_model */,
FC704C1B21D237FC00F98BAB /* vision_model */,
FC704C1621D2375300F98BAB /* superresoltion */,
);
name = models;
path = ../../models;
sourceTree = "<group>";
};
FC203FAB21CBFDBA00B37166 /* vision_model */ = {
FC704C1621D2375300F98BAB /* superresoltion */ = {
isa = PBXGroup;
children = (
FC203FAC21CBFDBA00B37166 /* mobilenet */,
FC203FAF21CBFDBA00B37166 /* yolo */,
FC704C1721D2375300F98BAB /* super_params */,
FC704C1821D2375300F98BAB /* super_model */,
);
path = superresoltion;
sourceTree = "<group>";
};
FC704C1B21D237FC00F98BAB /* vision_model */ = {
isa = PBXGroup;
children = (
FC704C1C21D237FC00F98BAB /* mobilenet */,
FC704C1F21D237FC00F98BAB /* yolo */,
);
path = vision_model;
sourceTree = "<group>";
};
FC203FAC21CBFDBA00B37166 /* mobilenet */ = {
FC704C1C21D237FC00F98BAB /* mobilenet */ = {
isa = PBXGroup;
children = (
FC203FAD21CBFDBA00B37166 /* combined_mobilenet_params */,
FC203FAE21CBFDBA00B37166 /* combined_mobilenet_model */,
FC704C1D21D237FC00F98BAB /* combined_mobilenet_params */,
FC704C1E21D237FC00F98BAB /* combined_mobilenet_model */,
);
path = mobilenet;
sourceTree = "<group>";
};
FC203FAF21CBFDBA00B37166 /* yolo */ = {
FC704C1F21D237FC00F98BAB /* yolo */ = {
isa = PBXGroup;
children = (
FC203FB021CBFDBA00B37166 /* yolo_params */,
FC203FB121CBFDBA00B37166 /* yolo_model */,
FC704C2021D237FC00F98BAB /* yolo_params */,
FC704C2121D237FC00F98BAB /* yolo_model */,
);
path = yolo;
sourceTree = "<group>";
......@@ -269,13 +283,15 @@
buildActionMask = 2147483647;
files = (
FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */,
FC203FB421CBFDBA00B37166 /* combined_mobilenet_model in Resources */,
FC203FB321CBFDBA00B37166 /* combined_mobilenet_params in Resources */,
FC704C2221D237FC00F98BAB /* combined_mobilenet_params in Resources */,
FC704C1921D2375300F98BAB /* super_params in Resources */,
FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */,
FC203FB521CBFDBA00B37166 /* yolo_params in Resources */,
FC203FB621CBFDBA00B37166 /* yolo_model in Resources */,
FC704C1A21D2375300F98BAB /* super_model in Resources */,
FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */,
FC203FB221CBFDBA00B37166 /* test.jpg in Resources */,
FC704C2321D237FC00F98BAB /* combined_mobilenet_model in Resources */,
FC704C2421D237FC00F98BAB /* yolo_params in Resources */,
FC704C2521D237FC00F98BAB /* yolo_model in Resources */,
);
runOnlyForDeploymentPostprocessing = 0;
};
......
......@@ -21,25 +21,32 @@ import MetalPerformanceShaders
var platform: Platform = .GPU
let threadSupport: [(Platform, String)] = [(.GPU, "GPU"), (.CPU, "CPU")]
let netSupport: [SupportModel : Net] = [
.super_resolution : SuperResolutionNet.init(device: MetalHelper.shared.device),
.yolo : YoloNet.init(device: MetalHelper.shared.device),
.mobilenet_combined : MobileNetCombined.init(device: MetalHelper.shared.device)]
//.mobilenet_ssd : Runner.init(inNet: MobileNet_ssd_hand.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform),
let modelHelperMap: [SupportModel : Runner] = [
.yolo : Runner.init(inNet: YoloNet.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform),
.mobilenet_combined : Runner.init(inNet: MobileNetCombined.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform)]
.super_resolution : Runner.init(inNet: netSupport[.super_resolution]!, commandQueue: MetalHelper.shared.queue, inPlatform: platform),
.yolo : Runner.init(inNet: YoloNet.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform),
.mobilenet_combined : Runner.init(inNet: MobileNetCombined.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform)]
//, .genet : Genet.init()
//let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()]
let netSupport: [SupportModel : Net] = [.yolo : YoloNet.init(device: MetalHelper.shared.device), .mobilenet_combined : MobileNetCombined.init(device: MetalHelper.shared.device)]
enum SupportModel: String{
// case mobilenet = "mobilenet"
// case mobilenet_ssd = "mobilenetssd"
case yolo = "yolo"
// case mobilenet_ssd = "mobilenetssd"
case yolo = "yolo"
case mobilenet_combined = "mobilenet_combined"
case super_resolution = "superresoltion"
static func supportedModels() -> [SupportModel] {
// .mobilenet,
// .mobilenet_ssd,
return [.yolo, .mobilenet_combined]
return [.super_resolution, .yolo, .mobilenet_combined]
}
}
......@@ -50,8 +57,8 @@ class ViewController: UIViewController {
@IBOutlet weak var modelPickerView: UIPickerView!
@IBOutlet weak var threadPickerView: UIPickerView!
@IBOutlet weak var videoView: UIView!
// var videoCapture: VideoCapture!
// var videoCapture: VideoCapture!
var selectImage: UIImage?
var inputPointer: UnsafeMutablePointer<Float32>?
var modelType: SupportModel = SupportModel.supportedModels()[0]
......@@ -62,18 +69,19 @@ class ViewController: UIViewController {
var threadNum = 1
@IBAction func loadAct(_ sender: Any) {
runner = Runner.init(inNet: netSupport[modelType]!, commandQueue: MetalHelper.shared.queue, inPlatform: platform)
runner = Runner.init(inNet: netSupport[modelType]!, commandQueue: MetalHelper.shared.queue, inPlatform: platform)
if platform == .CPU {
if inputPointer == nil {
inputPointer = runner.preproccess(image: selectImage!.cgImage!)
}
} else if platform == .GPU {
if self.toPredictTexture == nil {
runner.getTexture(image: selectImage!.cgImage!) {[weak self] (texture) in
runner.getTexture(image: selectImage!.cgImage!) { [weak self] (texture) in
self?.toPredictTexture = texture
}
}
} else {
fatalError( " unsupport " )
......@@ -106,27 +114,21 @@ class ViewController: UIViewController {
return
}
// for _ in 0..<1{
// runner.predict(texture: inTexture) { (success, resultHolder) in
// resultHolder?.releasePointer()
// }
// }
let startDate = Date.init()
for i in 0..<max {
runner.predict(texture: inTexture) { [weak self] (success, resultHolder) in
self.runner.predict(texture: inTexture) { [weak self] (success, resultHolder) in
guard let sSelf = self else {
fatalError()
}
if success {
if i == max - 1 {
let time = Date.init().timeIntervalSince(startDate)
print(Array<Any>.floatArrWithBuffer(floatArrBuffer: resultHolder!.result!, count: resultHolder!.capacity).strideArray())
DispatchQueue.main.async {
// print(resultHolder!.result![0])
// print(resultHolder!.result![0])
sSelf.resultTextView.text = sSelf.runner.net.resultStr(res: resultHolder!)
sSelf.elapsedTimeLabel.text = "平均耗时: \(time/Double(max) * 1000.0) ms"
}
}
}
......@@ -134,12 +136,9 @@ class ViewController: UIViewController {
DispatchQueue.main.async {
resultHolder?.releasePointer()
}
// print("释放")
}
// print("sleep before ")
// usleep(33000)
// print("sleep after ")
}
case .CPU:
guard let inInputPointer = inputPointer else {
fatalError( " need input pointer " )
......@@ -161,12 +160,15 @@ class ViewController: UIViewController {
if i == max - 1 {
let time = Date.init().timeIntervalSince(startDate)
DispatchQueue.main.async {
// sSelf.resultTextView.text = sSelf.runner.net.resultStr(res: res)
// sSelf.resultTextView.text = sSelf.runner.net.resultStr(res: res)
sSelf.elapsedTimeLabel.text = "平均耗时: \(time/Double(max) * 1000.0) ms"
}
}
}
res?.releaseOutput()
print(" predict done -- 123 ")
}
}
}
......@@ -180,36 +182,36 @@ class ViewController: UIViewController {
threadPickerView.delegate = self
threadPickerView.dataSource = self
if let image = UIImage.init(named: "test.jpg") {
selectImage = image
selectImageView.image = image
selectImage = image
selectImageView.image = image
} else {
print("请添加测试图片")
print("请添加测试图片")
}
// if platform == .CPU {
// inputPointer = runner.preproccess(image: selectImage!.cgImage!)
// } else if platform == .GPU {
// runner.getTexture(image: selectImage!.cgImage!) {[weak self] (texture) in
// self?.toPredictTexture = texture
// }
// } else {
// fatalError( " unsupport " )
// }
// if platform == .CPU {
// inputPointer = runner.preproccess(image: selectImage!.cgImage!)
// } else if platform == .GPU {
// runner.getTexture(image: selectImage!.cgImage!) {[weak self] (texture) in
// self?.toPredictTexture = texture
// }
// } else {
// fatalError( " unsupport " )
// }
// videoCapture = VideoCapture.init(device: MetalHelper.shared.device, orientation: .portrait, position: .back)
// videoCapture.fps = 30
// videoCapture.delegate = self
// videoCapture.setUp { (success) in
// DispatchQueue.main.async {
// if let preViewLayer = self.videoCapture.previewLayer {
// self.videoView.layer.addSublayer(preViewLayer)
// self.videoCapture.previewLayer?.frame = self.videoView.bounds
// }
// self.videoCapture.start()
// }
// }
// videoCapture = VideoCapture.init(device: MetalHelper.shared.device, orientation: .portrait, position: .back)
// videoCapture.fps = 30
// videoCapture.delegate = self
// videoCapture.setUp { (success) in
// DispatchQueue.main.async {
// if let preViewLayer = self.videoCapture.previewLayer {
// self.videoView.layer.addSublayer(preViewLayer)
// self.videoCapture.previewLayer?.frame = self.videoView.bounds
// }
// self.videoCapture.start()
// }
// }
}
}
......@@ -276,25 +278,11 @@ extension ViewController: VideoCaptureDelegate{
func predictTexture(texture: MTLTexture){
runner.scaleTexture(input: texture) { (scaledTexture) in
self.runner.predict(texture: scaledTexture, completion: { (success, resultHolder) in
// print(resultHolder!.result![0])
// print(resultHolder!.result![0])
resultHolder?.releasePointer()
})
}
}
// @available(iOS 10.0, *)
// func videoCapture(_ capture: VideoCapture, didCaptureVideoTexture texture: MTLTexture?, timestamp: CMTime) {
//// if !bool1 {
//// DispatchQueue.main.asyncAfter(deadline: DispatchTime.init(uptimeNanoseconds: 500000000)) {
// self.predictTexture(texture: texture!)
//// }
//
//
//// bool1 = true
//// }
//
// }
}
......
......@@ -64,6 +64,7 @@
FC0E2DBE20EE460D009C1FAC /* BatchNormKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC0E2DBD20EE460D009C1FAC /* BatchNormKernel.swift */; };
FC0E2DC020EE461F009C1FAC /* ElementwiseAddKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */; };
FC1B16B320EC9A4F00678B91 /* Kernels.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC1B16B220EC9A4F00678B91 /* Kernels.metal */; };
FC1CF3F721D4B4C400F7392E /* Runner.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC1CF3F621D4B4C400F7392E /* Runner.swift */; };
FC292C5421421B2F00CF622F /* PaddleMobileGPU.h in Headers */ = {isa = PBXBuildFile; fileRef = FC292C5321421B2E00CF622F /* PaddleMobileGPU.h */; settings = {ATTRIBUTES = (Public, ); }; };
FC292C5621421B4600CF622F /* PaddleMobileGPU.m in Sources */ = {isa = PBXBuildFile; fileRef = FC292C5521421B4600CF622F /* PaddleMobileGPU.m */; };
FC292C81214255BD00CF622F /* CPUCompute.mm in Sources */ = {isa = PBXBuildFile; fileRef = FC292C7C214255BC00CF622F /* CPUCompute.mm */; };
......@@ -74,11 +75,11 @@
FC3602CC2108819F00FACB58 /* PaddleMobileUnitTest.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */; };
FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74820F0B954007C0C6D /* ConvKernel.metal */; };
FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */; };
FC4FD9752140E1DE0073E130 /* PaddleMobile.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC4FD9742140E1DE0073E130 /* PaddleMobile.swift */; };
FC4FD9792140E4980073E130 /* PaddleMobileCPU.h in Headers */ = {isa = PBXBuildFile; fileRef = FC4FD9772140E4980073E130 /* PaddleMobileCPU.h */; settings = {ATTRIBUTES = (Public, ); }; };
FC4FD97A2140E4980073E130 /* libpaddle-mobile.a in Frameworks */ = {isa = PBXBuildFile; fileRef = FC4FD9782140E4980073E130 /* libpaddle-mobile.a */; };
FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */; };
FC60DB8920E9AAA500FF203F /* MetalExtension.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC60DB8820E9AAA500FF203F /* MetalExtension.swift */; };
FC704C2721D2385100F98BAB /* SuperResolutionNet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC704C2621D2385100F98BAB /* SuperResolutionNet.swift */; };
FC803BBF214CB65A0094B8E5 /* ConvAddPreluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC803BBE214CB65A0094B8E5 /* ConvAddPreluOp.swift */; };
FC803BC1214CB77A0094B8E5 /* ConvAddPreluKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC803BC0214CB77A0094B8E5 /* ConvAddPreluKernel.swift */; };
FC803BC3214CB79C0094B8E5 /* ConvAddPreluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC803BC2214CB79C0094B8E5 /* ConvAddPreluKernel.metal */; };
......@@ -87,6 +88,7 @@
FC803BC9214CFC8D0094B8E5 /* FetchKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC803BC8214CFC8D0094B8E5 /* FetchKernel.metal */; };
FC82735920E3C04200BE430A /* OpCreator.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC82735820E3C04200BE430A /* OpCreator.swift */; };
FC9A19E32148C31300CD9CBF /* MobilenetSSD_AR.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9A19E22148C31300CD9CBF /* MobilenetSSD_AR.swift */; };
FC9C2A0D21D3D185005856C6 /* FetchKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC9C2A0C21D3D185005856C6 /* FetchKernel.inc.metal */; };
FC9D037920E229E4000F735A /* OpParam.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037820E229E4000F735A /* OpParam.swift */; };
FC9D038020E22FBB000F735A /* FeedOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037F20E22FBB000F735A /* FeedOp.swift */; };
FC9D038220E2312E000F735A /* FetchOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D038120E2312E000F735A /* FetchOp.swift */; };
......@@ -203,6 +205,7 @@
FC0E2DBD20EE460D009C1FAC /* BatchNormKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = BatchNormKernel.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>"; };
FC1CF3F621D4B4C400F7392E /* Runner.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Runner.swift; sourceTree = "<group>"; };
FC292C5321421B2E00CF622F /* PaddleMobileGPU.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = PaddleMobileGPU.h; sourceTree = "<group>"; };
FC292C5521421B4600CF622F /* PaddleMobileGPU.m */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.objc; path = PaddleMobileGPU.m; sourceTree = "<group>"; };
FC292C7C214255BC00CF622F /* CPUCompute.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = CPUCompute.mm; sourceTree = "<group>"; };
......@@ -213,12 +216,12 @@
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>"; };
FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ProgramOptimize.swift; sourceTree = "<group>"; };
FC4FD9742140E1DE0073E130 /* PaddleMobile.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = PaddleMobile.swift; sourceTree = "<group>"; };
FC4FD9772140E4980073E130 /* PaddleMobileCPU.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = PaddleMobileCPU.h; sourceTree = "<group>"; };
FC4FD9782140E4980073E130 /* libpaddle-mobile.a */ = {isa = PBXFileReference; lastKnownFileType = archive.ar; path = "libpaddle-mobile.a"; sourceTree = "<group>"; };
FC4FD97D2140F2C30073E130 /* libstdc++.tbd */ = {isa = PBXFileReference; lastKnownFileType = "sourcecode.text-based-dylib-definition"; name = "libstdc++.tbd"; path = "usr/lib/libstdc++.tbd"; sourceTree = SDKROOT; };
FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Texture2DTo2DArrayKernel.swift; sourceTree = "<group>"; };
FC60DB8820E9AAA500FF203F /* MetalExtension.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MetalExtension.swift; sourceTree = "<group>"; };
FC704C2621D2385100F98BAB /* SuperResolutionNet.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = SuperResolutionNet.swift; sourceTree = "<group>"; };
FC803BBE214CB65A0094B8E5 /* ConvAddPreluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvAddPreluOp.swift; sourceTree = "<group>"; };
FC803BC0214CB77A0094B8E5 /* ConvAddPreluKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvAddPreluKernel.swift; sourceTree = "<group>"; };
FC803BC2214CB79C0094B8E5 /* ConvAddPreluKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvAddPreluKernel.metal; sourceTree = "<group>"; };
......@@ -227,6 +230,7 @@
FC803BC8214CFC8D0094B8E5 /* FetchKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = FetchKernel.metal; sourceTree = "<group>"; };
FC82735820E3C04200BE430A /* OpCreator.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpCreator.swift; sourceTree = "<group>"; };
FC9A19E22148C31300CD9CBF /* MobilenetSSD_AR.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobilenetSSD_AR.swift; sourceTree = "<group>"; };
FC9C2A0C21D3D185005856C6 /* FetchKernel.inc.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = FetchKernel.inc.metal; sourceTree = "<group>"; };
FC9D037820E229E4000F735A /* OpParam.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpParam.swift; sourceTree = "<group>"; };
FC9D037F20E22FBB000F735A /* FeedOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FeedOp.swift; sourceTree = "<group>"; };
FC9D038120E2312E000F735A /* FetchOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FetchOp.swift; sourceTree = "<group>"; };
......@@ -339,6 +343,7 @@
FC039B6C20E11C3C0081E9F8 /* paddle-mobile */ = {
isa = PBXGroup;
children = (
FC704C2621D2385100F98BAB /* SuperResolutionNet.swift */,
C28FDF8221B7858F0054EFAC /* MobileNetCombined.swift */,
C28FDF8321B7858F0054EFAC /* YoloNet.swift */,
FCE9D7B6214F869000B520C3 /* Net.swift */,
......@@ -351,13 +356,13 @@
FC292C5521421B4600CF622F /* PaddleMobileGPU.m */,
FC292C5321421B2E00CF622F /* PaddleMobileGPU.h */,
FC4FD9762140E4920073E130 /* CPU */,
FC4FD9742140E1DE0073E130 /* PaddleMobile.swift */,
FC039BAE20E11CC20081E9F8 /* Program */,
FC039BA320E11CBC0081E9F8 /* Operators */,
FC039B9C20E11CB20081E9F8 /* framework */,
FC039B9320E11C9A0081E9F8 /* Common */,
FC039B6D20E11C3C0081E9F8 /* paddle_mobile.h */,
FC039B6E20E11C3C0081E9F8 /* Info.plist */,
FC1CF3F621D4B4C400F7392E /* Runner.swift */,
);
path = "paddle-mobile";
sourceTree = "<group>";
......@@ -540,6 +545,7 @@
FCE9D7B8214FAA4800B520C3 /* NMSFetchResultKernel.metal */,
FCE3A1B02153E90F00C37CDE /* ElementwiseAddPreluKernel.inc.metal */,
FCE3A1B22153E91900C37CDE /* ElementwiseAddPreluKernel.metal */,
FC9C2A0C21D3D185005856C6 /* FetchKernel.inc.metal */,
);
path = metal;
sourceTree = "<group>";
......@@ -650,6 +656,7 @@
buildActionMask = 2147483647;
files = (
FC9D038020E22FBB000F735A /* FeedOp.swift in Sources */,
FC9C2A0D21D3D185005856C6 /* FetchKernel.inc.metal in Sources */,
4AA1EAAA214F53D800D0F791 /* BoxCoder.inc.metal in Sources */,
FC039B9F20E11CB20081E9F8 /* Tensor.swift in Sources */,
FC803BC9214CFC8D0094B8E5 /* FetchKernel.metal in Sources */,
......@@ -701,7 +708,6 @@
FC039B9920E11C9A0081E9F8 /* Types.swift in Sources */,
FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */,
FCA3A1632132A4AC00084FE5 /* ReshapeKernel.metal in Sources */,
FC4FD9752140E1DE0073E130 /* PaddleMobile.swift in Sources */,
FCBCCC592122F42700D94F7E /* ConvBNReluOp.swift in Sources */,
FC039BA920E11CBC0081E9F8 /* ConvOp.swift in Sources */,
FC9D038420E23B01000F735A /* Texture.swift in Sources */,
......@@ -715,6 +721,7 @@
FC039BBF20E11CC20081E9F8 /* Attribute.swift in Sources */,
4AA1EA8E2146647F00D0F791 /* SplitKernel.swift in Sources */,
FCD04E7420F3437E0007374F /* ConvAddKernel.swift in Sources */,
FC1CF3F721D4B4C400F7392E /* Runner.swift in Sources */,
FC039BB920E11CC20081E9F8 /* Scope.swift in Sources */,
FC292C5621421B4600CF622F /* PaddleMobileGPU.m in Sources */,
FCD04E6620F314C50007374F /* PoolOp.swift in Sources */,
......@@ -740,6 +747,7 @@
FCBCCC5D2122F8A100D94F7E /* DepthwiseConvOp.swift in Sources */,
FCE3A1AF2153E8EE00C37CDE /* ElementwiseAddPreluKernel.swift in Sources */,
FCE9D7B7214F869000B520C3 /* Net.swift in Sources */,
FC704C2721D2385100F98BAB /* SuperResolutionNet.swift in Sources */,
FC0E2DBE20EE460D009C1FAC /* BatchNormKernel.swift in Sources */,
FC039BAB20E11CBC0081E9F8 /* Operator.swift in Sources */,
C28FDF8421B7858F0054EFAC /* MobileNetCombined.swift in Sources */,
......
......@@ -110,8 +110,18 @@ extension Array {
return newArray
}
}
public static func floatArrWithBuffer(floatArrBuffer: UnsafeMutablePointer<Float32>, count: Int) -> [Float32] {
var arr: [Float32] = []
for i in 0..<count {
arr.append(floatArrBuffer[i])
}
return arr
}
}
extension String{
func cStr() -> UnsafePointer<Int8>? {
return (self as NSString).utf8String
......
......@@ -501,7 +501,7 @@ public extension MTLTexture {
} else {
fatalError(" 目前还不支持其他类型 ")
}
print(textureArray.count)
var output: [Float32] = []
for s in 0..<arrayLength {
for c in 0..<4{
......
......@@ -250,28 +250,29 @@ extension InputTexture: Variant {
}
extension MTLTexture where Self: Variant {
}
class FetchHolder: Variant {
var resultBuffer: MTLBuffer?
var dim: [Int]
var dim: Dim
var capacity: Int
var paddedCapacity: Int
init(inCapacity: Int, inDim: [Int]) {
capacity = inCapacity
init(inPaddedCapacity: Int, inDim: Dim) {
paddedCapacity = inPaddedCapacity
capacity = inDim.numel()
dim = inDim
}
func initBuffer(device: MTLDevice) {
resultBuffer = device.makeBuffer(length: capacity * 4, options: [])
resultBuffer = device.makeBuffer(length: paddedCapacity * 4, options: [])
}
var result: UnsafeMutablePointer<Float32> {
guard let inResultBuffer = resultBuffer else {
fatalError()
}
return inResultBuffer.contents().bindMemory(to: Float32.self, capacity: capacity)
return inResultBuffer.contents().bindMemory(to: Float32.self, capacity: paddedCapacity)
}
}
......
......@@ -24,7 +24,7 @@ public class Genet: Net {
paramPath = Bundle.main.path(forResource: "genet_params", ofType: nil) ?! "para null"
modelDir = ""
preprocessKernel = GenetPreProccess.init(device: device)
dim = (n: 1, h: 128, w: 128, c: 3)
inputDim_ = Dim.init(inDim: [1, 128, 128, 3])
}
@objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
......@@ -36,7 +36,7 @@ public class Genet: Net {
paramPath = ""
modelDir = ""
preprocessKernel = GenetPreProccess.init(device: device)
dim = (n: 1, h: 128, w: 128, c: 3)
inputDim_ = Dim.init(inDim: [1, 128, 128, 3])
}
class GenetPreProccess: CusomKernel {
......
......@@ -52,9 +52,7 @@ class MobileNet: Net{
}
return s.joined(separator: "\n")
}
override init(device: MTLDevice) {
super.init(device: device)
means = [123.68, 116.78, 103.94]
......@@ -64,7 +62,7 @@ class MobileNet: Net{
paramPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null"
modelDir = ""
preprocessKernel = MobilenetPreProccess.init(device: device)
dim = (n: 1, h: 224, w: 224, c: 3)
inputDim_ = Dim.init(inDim: [1, 224, 224, 3])
}
}
......@@ -9,40 +9,41 @@
import Foundation
public class MobileNetCombined: Net {
@objc public override init(device: MTLDevice) {
super.init(device: device)
means = [0, 0, 0]
scale = 1
except = 0
modelPath = Bundle.main.path(forResource: "combined_mobilenet_model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "combined_mobilenet_params", ofType: nil) ?! "para null"
modelDir = ""
//preprocessKernel = GenetPreProccess.init(device: device)
dim = (n: 1, h: 416, w: 416, c: 3)
}
@objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
super.init(device:device,paramPointer:paramPointer,paramSize:paramSize,modePointer:modePointer,modelSize:modelSize)
means = [0, 0, 0]
scale = 1
except = 0
modelPath = ""
paramPath = ""
modelDir = ""
//preprocessKernel = GenetPreProccess.init(device: device)
dim = (n: 1, h: 416, w: 416, c: 3)
}
// class GenetPreProccess: CusomKernel {
// init(device: MTLDevice) {
// let s = CusomKernel.Shape.init(inWidth: 128, inHeight: 128, inChannel: 3)
// super.init(device: device, inFunctionName: "genet_preprocess", outputDim: s, usePaddleMobileLib: false)
// }
// }
override public func resultStr(res: ResultHolder) -> String {
// fatalError()
return " \(res.result![0]) ... "
}
@objc public override init(device: MTLDevice) {
super.init(device: device)
means = [0, 0, 0]
scale = 1
except = 0
modelPath = Bundle.main.path(forResource: "combined_mobilenet_model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "combined_mobilenet_params", ofType: nil) ?! "para null"
modelDir = ""
//preprocessKernel = GenetPreProccess.init(device: device)
inputDim_ = Dim.init(inDim: [1, 416, 416, 3])
}
@objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
super.init(device:device,paramPointer:paramPointer,paramSize:paramSize,modePointer:modePointer,modelSize:modelSize)
means = [0, 0, 0]
scale = 1
except = 0
modelPath = ""
paramPath = ""
modelDir = ""
//preprocessKernel = GenetPreProccess.init(device: device)
inputDim_ = Dim.init(inDim: [1, 416, 416, 3])
}
// class GenetPreProccess: CusomKernel {
// init(device: MTLDevice) {
// let s = CusomKernel.Shape.init(inWidth: 128, inHeight: 128, inChannel: 3)
// super.init(device: device, inFunctionName: "genet_preprocess", outputDim: s, usePaddleMobileLib: false)
// }
// }
override public func resultStr(res: ResultHolder) -> String {
// fatalError()
return " \(res.result![0]) ... "
}
}
......@@ -24,7 +24,7 @@ public class MobileNet_ssd_hand: Net{
paramPath = Bundle.main.path(forResource: "ssd_hand_params", ofType: nil) ?! "para null"
modelDir = ""
preprocessKernel = MobilenetssdPreProccess.init(device: device)
dim = (n: 1, h: 300, w: 300, c: 3)
inputDim_ = Dim.init(inDim: [1, 300, 300, 3])
}
@objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
......@@ -36,7 +36,7 @@ public class MobileNet_ssd_hand: Net{
paramPath = ""
modelDir = ""
preprocessKernel = MobilenetssdPreProccess.init(device: device)
dim = (n: 1, h: 300, w: 300, c: 3)
inputDim_ = Dim.init(inDim: [1, 300, 300, 3])
}
class MobilenetssdPreProccess: CusomKernel {
......
......@@ -24,7 +24,7 @@ public class MobileNet_ssd_AR: Net{
paramPath = Bundle.main.path(forResource: "ar_params", ofType: nil) ?! "para null"
modelDir = ""
preprocessKernel = MobilenetssdPreProccess.init(device: device)
dim = (n: 1, h: 160, w: 160, c: 3)
inputDim_ = Dim.init(inDim: [1, 160, 160, 3])
}
@objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
......@@ -36,7 +36,7 @@ public class MobileNet_ssd_AR: Net{
paramPath = ""
modelDir = ""
preprocessKernel = MobilenetssdPreProccess.init(device: device)
dim = (n: 1, h: 160, w: 160, c: 3)
inputDim_ = Dim.init(inDim: [1, 160, 160, 3])
}
class MobilenetssdPreProccess: CusomKernel {
......@@ -96,57 +96,57 @@ public class MobileNet_ssd_AR: Net{
}
override func updateProgram(program: Program) {
for i in [56, 66, 76, 86, 93, 99] {
let opDesc = program.programDesc.blocks[0].ops[i]
let output = opDesc.outputs["Out"]!.first!
let v = program.scope[output]!
let originTexture = v as! Texture<Float32>
originTexture.tensorDim = Dim.init(inDim: [originTexture.tensorDim[1] / 7, originTexture.tensorDim[0] * 7])
originTexture.dim = Dim.init(inDim: [1, 1, originTexture.dim[3] / 7, originTexture.dim[2] * 7])
originTexture.padToFourDim = Dim.init(inDim: [1, 1, originTexture.padToFourDim[3] / 7, originTexture.padToFourDim[2] * 7])
program.scope[output] = originTexture
if i == 99 {
opDesc.attrs["axis"] = 0
} else {
opDesc.attrs["shape"] = originTexture.tensorDim.dims.map { Int32($0) }
}
}
for i in [58, 59, 88, 89, 95, 96, 68, 69, 78, 79] {
let opDesc = program.programDesc.blocks[0].ops[i]
let output = opDesc.outputs["Out"]!.first!
let v = program.scope[output]!
let originTexture = v as! Texture<Float32>
originTexture.tensorDim = Dim.init(inDim: [originTexture.tensorDim[1], originTexture.tensorDim[2]])
opDesc.attrs["shape"] = originTexture.tensorDim.dims.map { Int32($0) }
}
for i in [60, 101, 90, 97, 70, 80] {
let opDesc = program.programDesc.blocks[0].ops[i]
let output = opDesc.outputs["Out"]!.first!
let v = program.scope[output]!
let originTexture = v as! Texture<Float32>
originTexture.tensorDim = Dim.init(inDim: [originTexture.tensorDim[1], originTexture.tensorDim[2]])
opDesc.attrs["axis"] = (opDesc.attrs["axis"]! as! Int) - 1
}
for i in [102] {
let opDesc = program.programDesc.blocks[0].ops[i]
for output in opDesc.outputs["Out"]! {
let v = program.scope[output]!
let originTexture = v as! Texture<Float32>
originTexture.tensorDim = Dim.init(inDim: [originTexture.tensorDim[1], originTexture.tensorDim[2]])
}
opDesc.attrs["axis"] = (opDesc.attrs["axis"]! as! Int) - 1
print(" split axis \(opDesc.attrs["axis"])")
}
// for i in [56, 66, 76, 86, 93, 99] {
// let opDesc = program.programDesc.blocks[0].ops[i]
// let output = opDesc.outputs["Out"]!.first!
// let v = program.scope[output]!
// let originTexture = v as! Texture
// originTexture.tensorDim = Dim.init(inDim: [originTexture.tensorDim[1] / 7, originTexture.tensorDim[0] * 7])
//
// originTexture.dim = Dim.init(inDim: [1, 1, originTexture.dim[3] / 7, originTexture.dim[2] * 7])
//
// originTexture.padToFourDim = Dim.init(inDim: [1, 1, originTexture.padToFourDim[3] / 7, originTexture.padToFourDim[2] * 7])
//
// program.scope[output] = originTexture
//
// if i == 99 {
// opDesc.attrs["axis"] = 0
// } else {
// opDesc.attrs["shape"] = originTexture.tensorDim.dims.map { Int32($0) }
// }
// }
//
// for i in [58, 59, 88, 89, 95, 96, 68, 69, 78, 79] {
// let opDesc = program.programDesc.blocks[0].ops[i]
// let output = opDesc.outputs["Out"]!.first!
// let v = program.scope[output]!
//
//
//
// let originTexture = v as! Texture
// originTexture.tensorDim = Dim.init(inDim: [originTexture.tensorDim[1], originTexture.tensorDim[2]])
// opDesc.attrs["shape"] = originTexture.tensorDim.dims.map { Int32($0) }
// }
//
// for i in [60, 101, 90, 97, 70, 80] {
// let opDesc = program.programDesc.blocks[0].ops[i]
// let output = opDesc.outputs["Out"]!.first!
// let v = program.scope[output]!
// let originTexture = v as! Texture
// originTexture.tensorDim = Dim.init(inDim: [originTexture.tensorDim[1], originTexture.tensorDim[2]])
// opDesc.attrs["axis"] = (opDesc.attrs["axis"]! as! Int) - 1
// }
//
// for i in [102] {
// let opDesc = program.programDesc.blocks[0].ops[i]
// for output in opDesc.outputs["Out"]! {
// let v = program.scope[output]!
// let originTexture = v as! Texture
// originTexture.tensorDim = Dim.init(inDim: [originTexture.tensorDim[1], originTexture.tensorDim[2]])
// }
// opDesc.attrs["axis"] = (opDesc.attrs["axis"]! as! Int) - 1
// print(" split axis \(opDesc.attrs["axis"])")
// }
// 99
}
......
......@@ -35,7 +35,22 @@ public class Net: NSObject {
var except: Int = 0
var means: [Float] = []
var scale: Float = 0.0
var dim: (n: Int, h: Int, w: Int, c: Int) = (n: 0, h: 0, w: 0, c: 0)
var needUpdateProgram = true
public var inputDim: Dim {
get{
return inputDim_
}
set{
if inputDim_ != newValue {
needUpdateProgram = true
}
inputDim_ = newValue
}
}
var inputDim_: Dim = Dim.init(inDim: [])
var preprocessKernel: CusomKernel? = nil
var paramPointer: UnsafeMutableRawPointer? = nil
var paramSize: Int = 0
......@@ -44,15 +59,17 @@ public class Net: NSObject {
var modelPath: String = ""
var paramPath: String = ""
var modelDir: String = ""
let device: MTLDevice
@objc public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
self.paramPointer = paramPointer
self.paramSize = paramSize
self.modelPointer = modePointer
self.modelSize = modelSize
super.init()
self.paramPointer = paramPointer
self.paramSize = paramSize
self.modelPointer = modePointer
self.modelSize = modelSize
self.device = device
super.init()
}
public func resultStr(res: ResultHolder) -> String {
fatalError()
}
......@@ -62,6 +79,7 @@ public class Net: NSObject {
}
@objc public init(device: MTLDevice) {
self.device = device
super.init()
}
......
......@@ -34,8 +34,8 @@ class BatchNormParam<P: PrecisionType>: OpParam {
throw error
}
}
let input: Texture<P>
var output: Texture<P>
let input: Texture
var output: Texture
let bias: Tensor<P>
let mean: Tensor<P>
let scale: Tensor<P>
......
......@@ -30,8 +30,8 @@ class BilinearInterpParam<P: PrecisionType>: OpParam {
fatalError()
}
}
let input: Texture<P>
var output: Texture<P>
let input: Texture
var output: Texture
let out_h: Int
let out_w: Int
}
......
......@@ -37,10 +37,10 @@ class BoxcoderParam<P: PrecisionType>: OpParam {
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 priorBox: Texture
let priorBoxVar: Texture
let targetBox: Texture
var output: Texture
let codeType: String
let boxNormalized: Bool
}
......
......@@ -22,7 +22,7 @@ class ConcatParam<P: PrecisionType>: OpParam {
fatalError()
}
for x in xlist {
guard let variant = inScope[x], let v = variant as? Texture<P> else {
guard let variant = inScope[x], let v = variant as? Texture else {
fatalError()
}
if transpose.count == 0 {
......@@ -40,8 +40,8 @@ class ConcatParam<P: PrecisionType>: OpParam {
throw error
}
}
var input: [Texture<P>] = []
var output: Texture<P>
var input: [Texture] = []
var output: Texture
var transpose: [Int] = []
let axis: Int
}
......
......@@ -34,12 +34,12 @@ class ConvAddAddPreluParam<P: PrecisionType>: OpParam {
}
}
let input: Texture<P>
let input: Texture
let y: Tensor<P>
let filter: Tensor<P>
let mode: String
let alpha: Tensor<P>
var output: Texture<P>
var output: Texture
let stride: [Int32]
let paddings: [Int32]
let dilations: [Int32]
......
......@@ -40,7 +40,7 @@ class ConvAddBatchNormReluParam<P: PrecisionType>: OpParam {
}
}
let input: Texture<P>
let input: Texture
let variance: Tensor<P>
let bias: Tensor<P>
......@@ -52,7 +52,7 @@ class ConvAddBatchNormReluParam<P: PrecisionType>: OpParam {
var newScale: MTLBuffer?
var newBiase: MTLBuffer?
var output: Texture<P>
var output: Texture
let stride: [Int32]
let paddings: [Int32]
let dilations: [Int32]
......
......@@ -32,11 +32,11 @@ class ConvAddParam<P: PrecisionType>: OpParam {
}
}
let input: Texture<P>
let input: Texture
let y: Tensor<P>
let filter: Tensor<P>
var output: Texture<P>
var output: Texture
let stride: [Int32]
let paddings: [Int32]
let dilations: [Int32]
......
......@@ -33,12 +33,12 @@ class ConvAddPreluParam<P: PrecisionType>: OpParam {
}
}
let input: Texture<P>
let input: Texture
let y: Tensor<P>
let filter: Tensor<P>
let mode: String
let alpha: Tensor<P>
var output: Texture<P>
var output: Texture
let stride: [Int32]
let paddings: [Int32]
let dilations: [Int32]
......
......@@ -36,8 +36,7 @@ class ConvBNReluParam<P: PrecisionType>: OpParam {
}
}
let input: Texture<P>
let input: Texture
let variance: Tensor<P>
let bias: Tensor<P>
let mean: Tensor<P>
......@@ -47,7 +46,7 @@ class ConvBNReluParam<P: PrecisionType>: OpParam {
var newScale: MTLBuffer?
var newBiase: MTLBuffer?
var output: Texture<P>
var output: Texture
let stride: [Int32]
let paddings: [Int32]
let dilations: [Int32]
......
......@@ -31,9 +31,9 @@ class ConvParam<P: PrecisionType>: OpParam {
}
}
let input: Texture<P>
let input: Texture
let filter: Tensor<P>
var output: Texture<P>
var output: Texture
let stride: [Int32]
let paddings: [Int32]
let dilations: [Int32]
......
......@@ -55,9 +55,9 @@ class ElementwiseAddParam<P: PrecisionType>: OpParam {
}
}
var inputX: Texture<P>
var inputY: Texture<P>
var output: Texture<P>
var inputX: Texture
var inputY: Texture
var output: Texture
var axis: Int
}
......
......@@ -59,9 +59,9 @@ class ElementwiseAddPreluParam<P: PrecisionType>: OpParam {
let mode: String
let alpha: Tensor<P>
var inputX: Texture<P>
var inputY: Texture<P>
var output: Texture<P>
var inputX: Texture
var inputY: Texture
var output: Texture
var axis: Int
}
......
......@@ -17,7 +17,7 @@ import MetalKit
import CoreMedia
class FeedParam<P: PrecisionType>: OpParam{
var output: Texture<P>
var output: Texture
var input: InputTexture {
return scope.input() as! InputTexture
}
......
......@@ -17,13 +17,13 @@ import Metal
class FetchParam<P: PrecisionType>: OpParam{
var output: FetchHolder
let input: Texture<P>
let input: Texture
let scope: Scope
required init(opDesc: OpDesc, inScope: Scope) throws {
scope = inScope
do {
input = try FetchParam.inputX(inputs: opDesc.inputs, from: inScope)
output = FetchHolder.init(inCapacity: input.numel(), inDim: input.tensorDim.dims)
output = FetchHolder.init(inPaddedCapacity: input.elementCount(), inDim: input.tensorDim)
scope.setOutput(output: output)
} catch let error {
throw error
......@@ -51,17 +51,13 @@ class FetchKernel<P: PrecisionType>: Kernel, Computable {
if param.input.transpose == [0, 2, 3, 1] {
super.init(device: device, inFunctionName: "fetch_half")
} else {
// fatalError(" not support ")
super.init(device: device, inFunctionName: "fetch_placeholder_half")
print(" not support ")
fatalError(" not support ")
}
} else if computePrecision == .Float32 {
if param.input.transpose == [0, 2, 3, 1] {
super.init(device: device, inFunctionName: "fetch")
super.init(device: device, inFunctionName: "fetch_float")
} else {
print(" not support ")
super.init(device: device, inFunctionName: "fetch_placeholder")
// fatalError(" not support ")
fatalError(" not support ")
}
} else {
fatalError(" not support ")
......@@ -84,5 +80,11 @@ class FetchOp<P: PrecisionType>: Operator< FetchKernel<P>, FetchParam<P>>, Runab
throw error
}
}
func delogOutput() {
print("fetch output: ")
let resArr = Array<Any>.floatArrWithBuffer(floatArrBuffer: self.para.output.result, count: self.para.output.capacity)
print(resArr.strideArray())
}
}
......@@ -25,8 +25,8 @@ class FlattenParam<P: PrecisionType>: OpParam {
throw error
}
}
let input: Texture<P>
var output: Texture<P>
let input: Texture
var output: Texture
let axis: Int
}
......
/* 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. */
#ifdef P
#define CONCAT3_(a, b, c) a ## _ ## b ## _ ## c
#define CONCAT2_(a, b) a ## _ ## b
#define CONCAT2(a, b) a ## b
#define FUNC(m, n, q) CONCAT3_(m, n, q)
#define FUNC_T(m, n) CONCAT2_(m, n)
#define VECTOR(p, n) CONCAT2(p, n)
kernel void FUNC_T(fetch, P)(texture2d_array<P, access::read> inTexture [[texture(0)]],
device float *output [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
if (gid.x >= inTexture.get_width() ||
gid.y >= inTexture.get_height() ||
gid.z >= inTexture.get_array_size()) {
return;
}
int input_width = inTexture.get_width();
int input_height = inTexture.get_height();
const VECTOR(P, 4) input = inTexture.read(gid.xy, gid.z);
int output_to = 4 * input_width * input_height;
output[gid.z * output_to + 0 * input_width * input_height + gid.y * input_width + gid.x] = input.x;
output[gid.z * output_to + 1 * input_width * input_height + gid.y * input_width + gid.x] = input.y;
output[gid.z * output_to + 2 * input_width * input_height + gid.y * input_width + gid.x] = input.z;
output[gid.z * output_to + 3 * input_width * input_height + gid.y * input_width + gid.x] = input.w;
}
#endif
......@@ -15,47 +15,18 @@
#include <metal_stdlib>
using namespace metal;
kernel void fetch(texture2d_array<float, access::read> inTexture [[texture(0)]],
device float *output [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
if (gid.x >= inTexture.get_width() ||
gid.y >= inTexture.get_height() ||
gid.z >= inTexture.get_array_size()) {
return;
}
int input_width = inTexture.get_width();
int input_height = inTexture.get_height();
const float4 input = inTexture.read(gid.xy, gid.z);
int output_to = 4 * input_width * input_height;
output[gid.z * output_to + 0 * input_width * input_height + gid.y * input_width + gid.x] = input.x;
output[gid.z * output_to + 1 * input_width * input_height + gid.y * input_width + gid.x] = input.y;
output[gid.z * output_to + 2 * input_width * input_height + gid.y * input_width + gid.x] = input.z;
output[gid.z * output_to + 3 * input_width * input_height + gid.y * input_width + gid.x] = input.w;
}
#define P float
#include "FetchKernel.inc.metal"
#undef P
#define P half
#include "FetchKernel.inc.metal"
#undef P
kernel void fetch_half(texture2d_array<half, access::read> inTexture [[texture(0)]],
device float * output [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
if (gid.x >= inTexture.get_width() ||
gid.y >= inTexture.get_height() ||
gid.z >= inTexture.get_array_size()) {
return;
}
int input_width = inTexture.get_width();
int input_height = inTexture.get_height();
const half4 input = inTexture.read(gid.xy, gid.z);
int output_to = 4 * input_width * input_height;
output[gid.z * output_to + 0 * input_width * input_height + gid.y * input_width + gid.x] = input.x;
output[gid.z * output_to + 1 * input_width * input_height + gid.y * input_width + gid.x] = input.y;
output[gid.z * output_to + 2 * input_width * input_height + gid.y * input_width + gid.x] = input.z;
output[gid.z * output_to + 3 * input_width * input_height + gid.y * input_width + gid.x] = input.w;
}
kernel void fetch_placeholder(texture2d_array<float, access::read> inTexture [[texture(0)]],
device float *output [[buffer(0)]],
......@@ -64,8 +35,6 @@ kernel void fetch_placeholder(texture2d_array<float, access::read> inTexture [[t
}
kernel void fetch_placeholder_half(texture2d_array<half, access::read> inTexture [[texture(0)]],
device float *output [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
device float *output [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
}
......@@ -22,18 +22,18 @@ class MulticlassNMSParam<P: PrecisionType>: OpParam {
bboxes = try MulticlassNMSParam.getFirstTensor(key: "BBoxes", map: opDesc.inputs, from: inScope)
output = try MulticlassNMSParam.outputOut(outputs: opDesc.outputs, from: inScope)
middleOutput = FetchHolder.init(inCapacity: scores.tensorDim.numel(), inDim: scores.tensorDim.dims)
middleOutput = FetchHolder.init(inPaddedCapacity: scores.tensorDim.numel(), inDim: scores.tensorDim)
bboxOutput = FetchHolder.init(inCapacity: bboxes.tensorDim.numel(), inDim: bboxes.tensorDim.dims)
bboxOutput = FetchHolder.init(inPaddedCapacity: bboxes.tensorDim.numel(), inDim: bboxes.tensorDim)
} catch let error {
throw error
}
}
var bboxOutput: FetchHolder
var middleOutput: FetchHolder
let scores: Texture<P>
let bboxes: Texture<P>
var output: Texture<P>
let scores: Texture
let bboxes: Texture
var output: Texture
}
class MulticlassNMSOp<P: PrecisionType>: Operator<MulticlassNMSKernel<P>, MulticlassNMSParam<P>>, Runable, Creator, InferShaperable{
......
......@@ -32,8 +32,8 @@ class PoolParam<P: PrecisionType>: OpParam {
}
// let buffer = input.metalTexture.buffer.contents().assumingMemoryBound(to: P.self)
}
let input: Texture<P>
var output: Texture<P>
let input: Texture
var output: Texture
var ksize: [Int32]
var stride: [Int32]
var padding: [Int32]
......
......@@ -28,8 +28,8 @@ class PreluParam<P: PrecisionType>: OpParam {
}
let mode: String
let alpha: Tensor<P>
let input: Texture<P>
var output: Texture<P>
let input: Texture
var output: Texture
}
class PreluOp<P: PrecisionType>: Operator<PreluKernel<P>, PreluParam<P>>, Runable, Creator, InferShaperable{
......
......@@ -53,10 +53,10 @@ class PriorBoxParam<P: PrecisionType>: OpParam {
var stepH: Float32
let offset: Float32
let input: Texture<P>
let inputImage: Texture<P>
var output: Texture<P>
let outputVariances: Texture<P>
let input: Texture
let inputImage: Texture
var output: Texture
let outputVariances: Texture
}
class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P>>, Runable, Creator, InferShaperable{
......
......@@ -25,8 +25,8 @@ class ReluParam<P: PrecisionType>: OpParam {
throw error
}
}
let input: Texture<P>
var output: Texture<P>
let input: Texture
var output: Texture
}
class ReluOp<P: PrecisionType>: Operator<ReluKernel<P>, ReluParam<P>>, Runable, Creator, InferShaperable{
......@@ -48,9 +48,9 @@ class ReluOp<P: PrecisionType>: Operator<ReluKernel<P>, ReluParam<P>>, Runable,
func delogOutput() {
print(" \(type) output: ")
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 device = para.output.metalTexture!.device
let outputArray: [Float32] = device.texture2tensor(texture: para.output.metalTexture, dim: para.output.tensorDim.dims, transpose: para.output.transpose)
print(outputArray.strideArray())
// let device = para.output.metalTexture!.device
// let outputArray: [Float32] = device.texture2tensor(texture: para.output.metalTexture, dim: para.output.tensorDim.dims, transpose: para.output.transpose)
// print(outputArray.strideArray())
}
}
......
......@@ -48,9 +48,9 @@ class ReshapeParam<P: PrecisionType>: OpParam {
throw error
}
}
let input: Texture<P>
let input: Texture
let shape: [Int32]
var output: Texture<P>
var output: Texture
}
class ReshapeOp<P: PrecisionType>: Operator<ReshapeKernel<P>, ReshapeParam<P>>, Runable, Creator, InferShaperable{
......
......@@ -24,8 +24,8 @@ class ShapeParam<P: PrecisionType>: OpParam {
throw error
}
}
var output: Texture<P>
let input: Texture<P>
var output: Texture
let input: Texture
}
class ShapeOp<P: PrecisionType>: Operator<ShapeKernel<P>, ShapeParam<P>>, Runable, Creator, InferShaperable{
......
......@@ -32,8 +32,8 @@ class SoftmaxParam<P: PrecisionType>: OpParam {
throw error
}
}
let input: Texture<P>
var output: Texture<P>
let input: Texture
var output: Texture
}
class SoftmaxOp<P: PrecisionType>: Operator<SoftmaxKernel<P>, SoftmaxParam<P>>, Runable, Creator, InferShaperable{
......
......@@ -19,7 +19,7 @@ class SplitParam<P: PrecisionType>: OpParam {
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
input = try SplitParam.inputX(inputs: opDesc.inputs, from: inScope)
output = Texture<P>.init(device: input.metalTexture!.device, inDim: input.dim)
output = Texture.init(device: input.metalTexture!.device, inDim: input.dim)
axis = try SplitParam.getAttr(key: "axis", attrs: opDesc.attrs)
sections = try SplitParam.getAttr(key: "sections", attrs: opDesc.attrs)
if axis < 0 {
......@@ -29,7 +29,7 @@ class SplitParam<P: PrecisionType>: OpParam {
fatalError()
}
for out in outlist {
guard let variant = inScope[out], let v = variant as? Texture<P> else {
guard let variant = inScope[out], let v = variant as? Texture else {
fatalError()
}
outputList.append(v)
......@@ -41,9 +41,9 @@ class SplitParam<P: PrecisionType>: OpParam {
}
var axis: Int
let input: Texture<P>
var output: Texture<P>
var outputList: [Texture<P>] = []
let input: Texture
var output: Texture
var outputList: [Texture] = []
var sections: [Int32] = []
}
......
......@@ -26,8 +26,8 @@ class TransposeParam<P: PrecisionType>: OpParam {
throw error
}
}
let input: Texture<P>
var output: Texture<P>
let input: Texture
var output: Texture
let axis: [Int32]
}
......
......@@ -69,6 +69,7 @@
}
-(void)predict:(id<MTLTexture>)texture withCompletion:(void (^)(BOOL, NSArray<NSNumber *> *))completion {
[runner predictWithTexture:texture completion:^(BOOL success, ResultHolder * _Nullable result) {
NSMutableArray<NSNumber *> *resultArray = [NSMutableArray arrayWithCapacity:result.capacity];
for (int i = 0; i < result.capacity; ++i) {
......
/* 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. */
//
// Runner.swift
// paddle-mobile
//
// Created by liuRuiLong on 2018/12/27.
// Copyright © 2018 orange. All rights reserved.
//
import Metal
import MetalKit
import Foundation
......@@ -30,7 +23,6 @@ class ScaleKernel: CusomKernel {
fatalError(" unsupport ")
}
}
}
@objc public class Runner: NSObject {
......@@ -53,6 +45,10 @@ class ScaleKernel: CusomKernel {
* inPlatform: 需要使用的平台, GPU or CPU
*/
@objc public init(inNet: Net, commandQueue: MTLCommandQueue?, inPlatform: Platform) {
guard inNet.inputDim.cout() == 4 else {
fatalError(" input dim count must 4 ")
}
net = inNet
queue = commandQueue
device = queue?.device
......@@ -63,12 +59,13 @@ class ScaleKernel: CusomKernel {
if platform == .CPU {
cpuPaddleMobile = PaddleMobileCPU.init()
}
numel = net.dim.n * net.dim.c * net.dim.h * net.dim.w
numel = net.inputDim.numel()
meansNumber = net.means.map { NSNumber.init(value: $0) }
dimsNum = [NSNumber.init(value: net.dim.n),
NSNumber.init(value: net.dim.c),
NSNumber.init(value: net.dim.h),
NSNumber.init(value: net.dim.w)]
dimsNum = [NSNumber.init(value: net.inputDim[0]),
NSNumber.init(value: net.inputDim[3]),
NSNumber.init(value: net.inputDim[1]),
NSNumber.init(value: net.inputDim[2])]
}
/**
......@@ -82,11 +79,11 @@ class ScaleKernel: CusomKernel {
}
let loader = Loader<Float32>.init()
do {
// program = try loader.load(device: inDevice, paramPointer: net.paramPointer!, paramSize: net.paramSize,modePointer:net.modelPointer!,modelSize:net.modelSize)
// program = try loader.load(device: inDevice, paramPointer: net.paramPointer!, paramSize: net.paramSize,modePointer:net.modelPointer!,modelSize:net.modelSize)
program = try loader.load(device: inDevice, modelPath: net.modelPath, paraPath: net.paramPath)
net.updateProgram(program: program!)
executor = try Executor<Float32>.init(inDevice: inDevice, inQueue: inQueue, inProgram: program!)
net.updateProgram(program: program!)
} catch let error {
print(error)
return false
......@@ -98,7 +95,6 @@ class ScaleKernel: CusomKernel {
}
@objc public func predict(inputPointer: UnsafeMutablePointer<Float32>, completion: @escaping ( _ success: Bool, _ result: PaddleMobileCPUResult?) -> Void) {
guard let res = cpuPaddleMobile?.predictInput(inputPointer, dim: dimsNum) else {
completion(false, nil)
return
......@@ -112,14 +108,15 @@ class ScaleKernel: CusomKernel {
* ( _ success: Bool, _ time:TimeInterval, _ resultArray: [Float32]) -> Void : 回调闭包, 三个参数分别为: 是否成功, 预测耗时, 结果数组
*/
@objc public func predict(texture: MTLTexture, completion: @escaping ( _ success: Bool, _ result: ResultHolder?) -> Void) {
net.updateProgram(program: program!)
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
try self.executor?.predict(input: texture, dim: self.net.inputDim, completionHandle: { [weak self] (res) in
guard let SSelf = self else {
fatalError( " self nil " )
}
let result = SSelf.net.fetchResult(paddleMobileRes: res)
completion(true, result)
}, preProcessKernle: self.net.preprocessKernel, except: self.net.except)
}, preProcessKernle: self.net.preprocessKernel, except: self.net.except)
} catch let error {
print(error)
completion(false, nil)
......@@ -132,21 +129,21 @@ class ScaleKernel: CusomKernel {
* cgImage: 需要预测的图片
* ( _ success: Bool, _ time:TimeInterval, _ resultArray: [Float32]) -> Void : 回调闭包, 三个参数分别为: 是否成功, 预测耗时, 结果数组
*/
// @objc public func predict(cgImage: CGImage, completion: @escaping ( _ success: Bool, _ 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 {
// let input = preproccess(image: cgImage)
// predict(inputPointer: input, completion: completion)
// input.deinitialize(count: numel)
// input.deallocate()
// }
// }
// @objc public func predict(cgImage: CGImage, completion: @escaping ( _ success: Bool, _ 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 {
// let input = preproccess(image: cgImage)
// predict(inputPointer: input, completion: completion)
// input.deinitialize(count: numel)
// input.deallocate()
// }
// }
/*
* 清理内存, 调用此函数后, 不能再使用, 需重新 load
......@@ -164,12 +161,18 @@ class ScaleKernel: CusomKernel {
@objc public func preproccess(image: CGImage) -> UnsafeMutablePointer<Float> {
let output = UnsafeMutablePointer<Float>.allocate(capacity: numel)
let means = net.means.map { NSNumber.init(value: $0) }
let dims = [NSNumber.init(value: net.dim.n),
NSNumber.init(value: net.dim.c),
NSNumber.init(value: net.dim.h),
NSNumber.init(value: net.dim.w)]
cpuPaddleMobile?.preprocess(image, output: output, means: means, scale: net.scale, dim: dims)
return output
if net.inputDim.cout() == 4 {
let dims = [NSNumber.init(value: net.inputDim[0]),
NSNumber.init(value: net.inputDim[3]),
NSNumber.init(value: net.inputDim[1]),
NSNumber.init(value: net.inputDim[2])]
cpuPaddleMobile?.preprocess(image, output: output, means: means, scale: net.scale, dim: dims)
return output
}
fatalError()
}
/*
......@@ -190,7 +193,7 @@ class ScaleKernel: CusomKernel {
fatalError( " make buffer error" )
}
let scaleKernel = ScaleKernel.init(device: inDevice, shape: CusomKernel.Shape.init(inWidth: net.dim.w, inHeight: net.dim.h, inChannel: 3))
let scaleKernel = ScaleKernel.init(device: inDevice, shape: CusomKernel.Shape.init(inWidth: net.inputDim[2], inHeight: net.inputDim[1], inChannel: 3))
do {
try scaleKernel.compute(inputTexuture: input, commandBuffer: buffer)
......@@ -205,5 +208,3 @@ class ScaleKernel: CusomKernel {
buffer.commit()
}
}
/* 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
public class SuperResolutionNet: Net{
override public func resultStr(res: ResultHolder) -> String {
return "未实现"
}
override public init(device: MTLDevice) {
super.init(device: device)
means = [0.0, 0.0, 0.0]
scale = 1.0
except = 0
modelPath = Bundle.main.path(forResource: "super_model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "super_params", ofType: nil) ?! "para null"
modelDir = ""
preprocessKernel = nil
inputDim_ = Dim.init(inDim: [1, Int(552 * 1.414), Int(310 * 1.414), 1])
}
override func updateProgram(program: Program) {
guard needUpdateProgram else {
return
}
// n h w c
for block in program.programDesc.blocks {
for varDesc in block.vars {
if !varDesc.persistable {
if varDesc.type == .LodTensor {
let varEle = program.scope.vars[varDesc.name]
if let texture = varEle as? Texture {
let newDim = Dim.init(inDim: [texture.dim[0], inputDim[1], inputDim[2], texture.tensorDim[1]])
print(" var desc name " + varDesc.name + " new dim" + "\(newDim)")
texture.updateDims(inTensorDim: Dim.init(inDim: [texture.tensorDim[0], texture.tensorDim[1], inputDim[1], inputDim[2]]), inDim: newDim)
texture.initTexture(device: device, inTranspose: [0, 1, 2, 3], computePrecision: computePrecision)
let output: FetchHolder = program.scope.output() as! FetchHolder
output.dim = newDim
output.capacity = newDim.numel()
output.paddedCapacity = newDim.numel() * 4
output.initBuffer(device: device)
}
}
}
}
}
needUpdateProgram = false
}
}
......@@ -10,40 +10,41 @@ import Foundation
import Metal
public class YoloNet: Net {
@objc public override init(device: MTLDevice) {
super.init(device: device)
means = [0, 0, 0]
scale = 1
except = 0
modelPath = Bundle.main.path(forResource: "yolo_model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "yolo_params", ofType: nil) ?! "para null"
modelDir = ""
//preprocessKernel = GenetPreProccess.init(device: device)
dim = (n: 1, h: 224, w: 224, c: 3)
}
@objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
super.init(device:device,paramPointer:paramPointer,paramSize:paramSize,modePointer:modePointer,modelSize:modelSize)
means = [0, 0, 0]
scale = 1
except = 0
modelPath = ""
paramPath = ""
modelDir = ""
//preprocessKernel = GenetPreProccess.init(device: device)
dim = (n: 1, h: 224, w: 224, c: 3)
}
// class GenetPreProccess: CusomKernel {
// init(device: MTLDevice) {
// let s = CusomKernel.Shape.init(inWidth: 128, inHeight: 128, inChannel: 3)
// super.init(device: device, inFunctionName: "genet_preprocess", outputDim: s, usePaddleMobileLib: false)
// }
// }
override public func resultStr(res: ResultHolder) -> String {
// fatalError()
return " \(res.result![0]) ... "
}
@objc public override init(device: MTLDevice) {
super.init(device: device)
means = [0, 0, 0]
scale = 1
except = 0
modelPath = Bundle.main.path(forResource: "yolo_model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "yolo_params", ofType: nil) ?! "para null"
modelDir = ""
//preprocessKernel = GenetPreProccess.init(device: device)
inputDim_ = Dim.init(inDim: [1, 416, 416, 3])
}
@objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
super.init(device:device,paramPointer:paramPointer,paramSize:paramSize,modePointer:modePointer,modelSize:modelSize)
means = [0, 0, 0]
scale = 1
except = 0
modelPath = ""
paramPath = ""
modelDir = ""
//preprocessKernel = GenetPreProccess.init(device: device)
inputDim_ = Dim.init(inDim: [1, 416, 416, 3])
}
// class GenetPreProccess: CusomKernel {
// init(device: MTLDevice) {
// let s = CusomKernel.Shape.init(inWidth: 128, inHeight: 128, inChannel: 3)
// super.init(device: device, inFunctionName: "genet_preprocess", outputDim: s, usePaddleMobileLib: false)
// }
// }
override public func resultStr(res: ResultHolder) -> String {
// fatalError()
return " \(res.result![0]) ... "
}
}
......@@ -18,7 +18,11 @@ public struct Dim {
public init(inDim: [Int]) {
dims = inDim
}
public init(inDim: (n: Int, h: Int, w: Int, c: Int)) {
dims = [inDim.n, inDim.h, inDim.w, inDim.c]
}
mutating func swapeDimAt(index1: Int, index2: Int) {
dims.swapAt(index1, index2)
}
......@@ -34,6 +38,10 @@ public struct Dim {
public static func ==(left: Dim, right: Dim) -> Bool {
return left.dims == right.dims;
}
public static func !=(left: Dim, right: Dim) -> Bool {
return left.dims != right.dims;
}
public subscript(index: Int) -> Int {
return dims[index];
......
......@@ -15,7 +15,7 @@
import Foundation
let testTo = 81
let testTo = 22
var isTest = false
......@@ -44,19 +44,19 @@ public class GPUResultHolder {
extension GPUResultHolder: CustomDebugStringConvertible, CustomStringConvertible {
public var debugDescription: String {
// var str = ""
// str += "Dim: \(dim) \n value:[ "
// if resultArr.count < 20 {
// for d in resultArr {
// str += " \(d) "
// }
// } else {
// for d in stride(from: 0, to: resultArr.count, by: resultArr.count/20) {
// str += " \(resultArr[d]) "
// }
// }
// str += " ]"
// return str
// var str = ""
// str += "Dim: \(dim) \n value:[ "
// if resultArr.count < 20 {
// for d in resultArr {
// str += " \(d) "
// }
// } else {
// for d in stride(from: 0, to: resultArr.count, by: resultArr.count/20) {
// str += " \(resultArr[d]) "
// }
// }
// str += " ]"
// return str
fatalError()
}
......@@ -67,17 +67,19 @@ extension GPUResultHolder: CustomDebugStringConvertible, CustomStringConvertible
public class Executor<P: PrecisionType> {
var ops: [Runable & InferShaperable] = []
var preInputDim: Dim = Dim.init(inDim: [])
let program: Program
let device: MTLDevice
let inflightSemaphore: DispatchSemaphore
let queue: MTLCommandQueue
public init(inDevice:MTLDevice, inQueue: MTLCommandQueue, inProgram: Program) throws {
self.inflightSemaphore = DispatchSemaphore(value: 3)
self.inflightSemaphore = DispatchSemaphore(value: 1)
program = inProgram
device = inDevice
queue = inQueue
// print("before for ")
//print(program.scope.vars["fea_pyramid1_mbox_conf_flat.Flatten.output.1.tmp_0"])
// print("before for ")
//print(program.scope.vars["fea_pyramid1_mbox_conf_flat.Flatten.output.1.tmp_0"])
for block in inProgram.programDesc.blocks {
......@@ -85,13 +87,15 @@ public class Executor<P: PrecisionType> {
for i in 0..<block.ops.count {
let opDesc = block.ops[i]
do {
// print("in for i \(i): ")
// print(program.scope.vars["fea_pyramid1_mbox_conf_flat.Flatten.output.1.tmp_0"])
//
// if i == 56 {
// print(program.scope.vars["fea_pyramid1_mbox_conf_flat.Flatten.output.1.tmp_0"])
//
// }
// print("in for i \(i): ")
// print(program.scope.vars["fea_pyramid1_mbox_conf_flat.Flatten.output.1.tmp_0"])
//
// if i == 56 {
// print(program.scope.vars["fea_pyramid1_mbox_conf_flat.Flatten.output.1.tmp_0"])
//
// }
let op = try OpCreator<P>.shared.creat(device: inDevice, opDesc: opDesc, scope: inProgram.scope)
ops.append(op)
......@@ -102,11 +106,12 @@ public class Executor<P: PrecisionType> {
}
}
public func predict(input: MTLTexture, dim: [Int], completionHandle: @escaping (GPUResultHolder) -> Void, preProcessKernle: CusomKernel? = nil, except: Int = 0) throws {
public func predict(input: MTLTexture, dim: Dim, completionHandle: @escaping (GPUResultHolder) -> Void, preProcessKernle: CusomKernel? = nil, except: Int = 0) throws {
inflightSemaphore.wait()
guard let buffer = queue.makeCommandBuffer() else {
throw PaddleMobileError.predictError(message: "CommandBuffer is nil")
}
inflightSemaphore.wait()
let resInput: MTLTexture
if let inPre = preProcessKernle {
......@@ -121,7 +126,7 @@ public class Executor<P: PrecisionType> {
}
let beforeDate = Date.init()
let inputTexture = InputTexture.init(inMTLTexture: resInput, inExpectDim: Dim.init(inDim: dim))
let inputTexture = InputTexture.init(inMTLTexture: resInput, inExpectDim: dim)
program.scope.setInput(input: inputTexture)
//(ops.count - except)
for i in 0..<(ops.count - except) {
......@@ -140,39 +145,25 @@ public class Executor<P: PrecisionType> {
}
buffer.addCompletedHandler { [weak self] (commandbuffer) in
// let inputArr = resInput.toTensor(dim: (n: dim[0], c: dim[3], h: dim[1], w: dim[2]))
// print(inputArr.strideArray())
//
//// print(dim)
// writeToLibrary(fileName: "test_image_ssd_ar", array: inputArr)
// print(" write done ")
// print("write to library done")
// return
// print(inputArr)
//
// let stridableInput: [(index: Int, value: Float)] = input.stridableFloatArray()
// print(stridableInput)
//
// let _: Flo? = input.logDesc(header: "input: ", stridable: true)
// for i in 0..<self!.ops.count {
// let op = self!.ops[i]
// print(" 第 \(i) 个 op: ")
// op.delogOutput()
// }
// return;
// self!.ops[testTo - 2].delogOutput()
// self!.ops[testTo - 1].delogOutput()
// self!.ops[5].delogOutput()
// return
guard let SSelf = self else {
// return
fatalError()
}
//将输入写进文件
/*
let inputArr = resInput.toTensor(dim: (n: dim[0], c: dim[3], h: dim[1], w: dim[2]))
print(dim)
writeToLibrary(fileName: "test_image_ssd_ar", array: inputArr)
print(" write done ")
return
*/
/* 输出 op 计算结果
for op in SSelf.ops {
op.delogOutput()
}
*/
let afterDate = Date.init()
var resultHolder: GPUResultHolder
if except > 0 {
......@@ -180,17 +171,13 @@ public class Executor<P: PrecisionType> {
} else {
let outputVar: Variant = SSelf.program.scope.output()!
let output: FetchHolder = outputVar as! FetchHolder
// let beforeToTensorDate = Date.init()
resultHolder = GPUResultHolder.init(inDim: output.dim, inPointer: output.result, inCapacity: output.capacity, inElapsedTime: afterDate.timeIntervalSince(beforeDate))
// let timeToTensor = Date.init().timeIntervalSince(beforeToTensorDate)
// print(timeToTensor)
resultHolder = GPUResultHolder.init(inDim: output.dim.dims, inPointer: output.result, inCapacity: output.capacity, inElapsedTime: afterDate.timeIntervalSince(beforeDate))
}
completionHandle(resultHolder)
SSelf.inflightSemaphore.signal()
}
buffer.commit()
}
......
......@@ -150,6 +150,9 @@ public class Loader<P: PrecisionType> {
let originProgramDesc = ProgramDesc.init(protoProgram: protoProgram)
let programDesc = ProgramOptimize<P>.init().optimize(originProgramDesc: originProgramDesc)
// let programDesc = ProgramDesc.init(protoProgram: protoProgram)
print(programDesc)
guard programDesc.blocks.count > 0 else {
......@@ -210,7 +213,7 @@ public class Loader<P: PrecisionType> {
scope[varDesc.name] = tensor
} else {
let dim = Dim.init(inDim: tensorDesc.dims)
scope[varDesc.name] = Texture<P>.init(device: device, inDim: dim)
scope[varDesc.name] = Texture.init(device: device, inDim: dim)
}
} else {
if varDesc.name == fetchKey {
......
......@@ -70,7 +70,7 @@ extension InputTexture {
*/
public class Texture<P: PrecisionType>: Tensorial {
public class Texture: Tensorial {
var dim: Dim
public var tensorDim: Dim
public var padToFourDim: Dim
......@@ -78,6 +78,10 @@ public class Texture<P: PrecisionType>: Tensorial {
public var metalTexture: MTLTexture!
var transpose: [Int] = [0, 1, 2, 3]
func elementCount() -> Int {
return metalTexture.width * metalTexture.height * metalTexture.arrayLength * 4
}
func toTensor() -> [Float32] {
guard padToFourDim.cout() == 4 else {
fatalError("- not support -")
......@@ -92,6 +96,8 @@ public class Texture<P: PrecisionType>: Tensorial {
return metalTexture.realNHWC(dim: (n: padToFourDim[0], h: padToFourDim[1], w: padToFourDim[2], c: padToFourDim[3]))
}
func initTexture(device: MTLDevice, inTranspose: [Int] = [0, 1, 2, 3], computePrecision: ComputePrecision = .Float16) {
transpose = inTranspose
for i in 0..<(4 - tensorDim.cout()) {
......@@ -99,8 +105,8 @@ public class Texture<P: PrecisionType>: Tensorial {
fatalError()
}
}
let newDim = transpose.map { padToFourDim[$0] }
let newDim = transpose.map { padToFourDim[$0] }
let newLayout = transpose.map { layout.layoutWithDim[$0] }
layout = DataLayout.init(newLayout)
......@@ -139,6 +145,26 @@ public class Texture<P: PrecisionType>: Tensorial {
metalTexture = device.makeTexture(descriptor: tmpTextureDes) ?! " texture nil "
}
func updateDims(inTensorDim: Dim, inDim: Dim) {
var fourDim: Dim
if inDim.cout() == 4 {
fourDim = inDim
} else if inDim.cout() < 4 {
var fourDimNum: [Int] = []
for _ in 0..<(4 - inDim.cout()) {
fourDimNum.append(1)
}
fourDimNum.append(contentsOf: inDim.dims)
fourDim = Dim.init(inDim: fourDimNum)
} else {
fatalError(" not support ")
}
tensorDim = inTensorDim
dim = fourDim
padToFourDim = fourDim
}
init(device: MTLDevice, inDim: Dim) {
var fourDim: Dim
if inDim.cout() == 4 {
......
......@@ -62,6 +62,7 @@ static const char *g_imgfssd_ar = "../images/test_image_ssd_ar";
static const char *g_imgfssd_ar1 = "../images/003_0001.txt";
static const char *g_img = "../images/img.bin";
static const char *g_yolo_img = "../images/in_put_1_3_416_416_2";
static const char *g_super_img = "../images/test_image_super";
static const char *g_mobilenet_img = "../images/image";
using paddle_mobile::framework::DDim;
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册