Skip to content

P
Paddle-Lite

PaddlePaddle / Paddle-Lite

通知 331
Star 4
Fork 1
P
Paddle-Lite
提交 a789f668 编写于 作者: R Ruilong Liu 提交者： GitHub
浏览文件
操作

Merge pull request #837 from codeWorm2015/metal 

align result
上级 10944966 ec42180c
显示空白变更内容
内联 并排
Showing with 2517 addition and 2130 deletion
metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj
浏览文件 @ a789f668
......@@ -16,6 +16,7 @@
FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */; };
FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */ = {isa = PBXBuildFile; fileRef = FC918190211DBC3500B6F354 /* paddle-mobile.png */; };
FC918193211DC70500B6F354 /* iphone.JPG in Resources */ = {isa = PBXBuildFile; fileRef = FC918192211DC70500B6F354 /* iphone.JPG */; };
FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FCA3A16021313E1F00084FE5 /* hand.jpg */; };
FCBCCC522122EEDC00D94F7E /* ssd_hand_params in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC502122EEDC00D94F7E /* ssd_hand_params */; };
FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC512122EEDC00D94F7E /* ssd_hand_model */; };
FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC542122EF5400D94F7E /* MetalHelper.swift */; };
......@@ -56,6 +57,7 @@
FC039B8D20E11C560081E9F8 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; 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>"; };
FCA3A16021313E1F00084FE5 /* hand.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = hand.jpg; sourceTree = "<group>"; };
FCBCCC502122EEDC00D94F7E /* ssd_hand_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_params; sourceTree = "<group>"; };
FCBCCC512122EEDC00D94F7E /* ssd_hand_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_model; sourceTree = "<group>"; };
FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = "<group>"; };
......@@ -137,6 +139,7 @@
FC0E2C1D20EDC030009C1FAC /* images */ = {
isa = PBXGroup;
children = (
FCA3A16021313E1F00084FE5 /* hand.jpg */,
FC918192211DC70500B6F354 /* iphone.JPG */,
FC918190211DBC3500B6F354 /* paddle-mobile.png */,
FCDFD41A211D91C7005AB38B /* synset.txt */,
......@@ -245,6 +248,7 @@
FCDFD41B211D91C7005AB38B /* synset.txt in Resources */,
FCD04E6420F3146B0007374F /* model in Resources */,
FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */,
FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */,
FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */,
);
runOnlyForDeploymentPostprocessing = 0;
......
metal/paddle-mobile-demo/paddle-mobile-demo/Base.lproj/Main.storyboard
浏览文件 @ a789f668
......@@ -19,10 +19,10 @@
<rect key="frame" x="0.0" y="0.0" width="375" height="667"/>
<autoresizingMask key="autoresizingMask" widthSizable="YES" heightSizable="YES"/>
<subviews>
<imageView userInteractionEnabled="NO" contentMode="scaleAspectFit" horizontalHuggingPriority="251" verticalHuggingPriority="251" translatesAutoresizingMaskIntoConstraints="NO" id="ZZh-fw-LwK">
<imageView userInteractionEnabled="NO" contentMode="scaleAspectFit" horizontalHuggingPriority="251" verticalHuggingPriority="251" ambiguous="YES" image="hand.jpg" translatesAutoresizingMaskIntoConstraints="NO" id="ZZh-fw-LwK">
<rect key="frame" x="0.0" y="20" width="375" height="247"/>
</imageView>
<label opaque="NO" userInteractionEnabled="NO" contentMode="left" horizontalHuggingPriority="251" verticalHuggingPriority="251" text="Thread:" textAlignment="natural" lineBreakMode="tailTruncation" baselineAdjustment="alignBaselines" adjustsFontSizeToFit="NO" translatesAutoresizingMaskIntoConstraints="NO" id="2EB-m2-a3L">
<label opaque="NO" userInteractionEnabled="NO" contentMode="left" horizontalHuggingPriority="251" verticalHuggingPriority="251" ambiguous="YES" text="Thread:" textAlignment="natural" lineBreakMode="tailTruncation" baselineAdjustment="alignBaselines" adjustsFontSizeToFit="NO" translatesAutoresizingMaskIntoConstraints="NO" id="2EB-m2-a3L">
<rect key="frame" x="10" y="538" width="68" height="24"/>
<constraints>
<constraint firstAttribute="width" constant="68" id="Q5J-tq-JSX"/>
......@@ -32,19 +32,19 @@
<nil key="textColor"/>
<nil key="highlightedColor"/>
</label>
<pickerView contentMode="scaleToFill" translatesAutoresizingMaskIntoConstraints="NO" id="DlO-dk-RMr">
<pickerView contentMode="scaleToFill" ambiguous="YES" translatesAutoresizingMaskIntoConstraints="NO" id="DlO-dk-RMr">
<rect key="frame" x="88" y="510.5" width="287" height="80"/>
<constraints>
<constraint firstAttribute="height" constant="80" id="Sbi-05-Mwd"/>
</constraints>
</pickerView>
<pickerView contentMode="scaleToFill" translatesAutoresizingMaskIntoConstraints="NO" id="6MG-gv-hD5">
<pickerView contentMode="scaleToFill" ambiguous="YES" translatesAutoresizingMaskIntoConstraints="NO" id="6MG-gv-hD5">
<rect key="frame" x="85" y="401" width="290" height="80"/>
<constraints>
<constraint firstAttribute="height" constant="80" id="yAL-JY-G6b"/>
</constraints>
</pickerView>
<label opaque="NO" userInteractionEnabled="NO" contentMode="left" horizontalHuggingPriority="251" verticalHuggingPriority="251" text="Models" textAlignment="natural" lineBreakMode="tailTruncation" baselineAdjustment="alignBaselines" adjustsFontSizeToFit="NO" translatesAutoresizingMaskIntoConstraints="NO" id="avL-VK-Kha">
<label opaque="NO" userInteractionEnabled="NO" contentMode="left" horizontalHuggingPriority="251" verticalHuggingPriority="251" ambiguous="YES" text="Models" textAlignment="natural" lineBreakMode="tailTruncation" baselineAdjustment="alignBaselines" adjustsFontSizeToFit="NO" translatesAutoresizingMaskIntoConstraints="NO" id="avL-VK-Kha">
<rect key="frame" x="10" y="429" width="65" height="24"/>
<constraints>
<constraint firstAttribute="width" constant="65" id="6oA-g2-Xq4"/>
......@@ -54,7 +54,7 @@
<nil key="textColor"/>
<nil key="highlightedColor"/>
</label>
<button opaque="NO" contentMode="scaleToFill" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" showsTouchWhenHighlighted="YES" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="wUL-9N-u1V">
<button opaque="NO" contentMode="scaleToFill" ambiguous="YES" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" showsTouchWhenHighlighted="YES" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="wUL-9N-u1V">
<rect key="frame" x="16" y="597" width="63.5" height="30"/>
<color key="backgroundColor" white="0.0" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
<state key="normal" title="Image">
......@@ -64,7 +64,7 @@
<action selector="selectImageAct:" destination="BYZ-38-t0r" eventType="touchUpInside" id="5uR-SM-fKO"/>
</connections>
</button>
<button opaque="NO" contentMode="scaleToFill" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" showsTouchWhenHighlighted="YES" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="XpL-9M-UOp">
<button opaque="NO" contentMode="scaleToFill" ambiguous="YES" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" showsTouchWhenHighlighted="YES" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="XpL-9M-UOp">
<rect key="frame" x="109.5" y="597" width="63" height="30"/>
<color key="backgroundColor" white="0.0" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
<state key="normal" title="Load">
......@@ -74,7 +74,7 @@
<action selector="loadAct:" destination="BYZ-38-t0r" eventType="touchUpInside" id="fZ5-CQ-jCY"/>
</connections>
</button>
<button opaque="NO" contentMode="scaleToFill" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" showsTouchWhenHighlighted="YES" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="R90-Yf-S6g">
<button opaque="NO" contentMode="scaleToFill" ambiguous="YES" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" showsTouchWhenHighlighted="YES" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="R90-Yf-S6g">
<rect key="frame" x="202.5" y="597" width="63.5" height="30"/>
<color key="backgroundColor" white="0.0" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
<state key="normal" title="Predict">
......@@ -84,7 +84,7 @@
<action selector="predictAct:" destination="BYZ-38-t0r" eventType="touchUpInside" id="Iyy-sY-gt4"/>
</connections>
</button>
<button opaque="NO" contentMode="scaleToFill" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" showsTouchWhenHighlighted="YES" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="a3K-ri-NVs">
<button opaque="NO" contentMode="scaleToFill" ambiguous="YES" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" showsTouchWhenHighlighted="YES" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="a3K-ri-NVs">
<rect key="frame" x="296" y="597" width="63" height="30"/>
<color key="backgroundColor" white="0.0" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
<state key="normal" title="Clear">
......@@ -94,7 +94,7 @@
<action selector="clearAct:" destination="BYZ-38-t0r" eventType="touchUpInside" id="JYf-UX-rCR"/>
</connections>
</button>
<view contentMode="scaleToFill" translatesAutoresizingMaskIntoConstraints="NO" id="w7H-Sk-Rai">
<view contentMode="scaleToFill" ambiguous="YES" translatesAutoresizingMaskIntoConstraints="NO" id="w7H-Sk-Rai">
<rect key="frame" x="79.5" y="597" width="30" height="30"/>
<color key="backgroundColor" white="1" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
<constraints>
......@@ -102,7 +102,7 @@
<constraint firstAttribute="width" constant="30" id="vYd-Fc-KAj"/>
</constraints>
</view>
<view contentMode="scaleToFill" translatesAutoresizingMaskIntoConstraints="NO" id="T4O-nx-ciH">
<view contentMode="scaleToFill" ambiguous="YES" translatesAutoresizingMaskIntoConstraints="NO" id="T4O-nx-ciH">
<rect key="frame" x="266" y="597" width="30" height="30"/>
<color key="backgroundColor" white="1" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
<constraints>
......@@ -110,7 +110,7 @@
<constraint firstAttribute="width" constant="30" id="fXE-S7-ZXL"/>
</constraints>
</view>
<view contentMode="scaleToFill" translatesAutoresizingMaskIntoConstraints="NO" id="976-fk-Kx2">
<view contentMode="scaleToFill" ambiguous="YES" translatesAutoresizingMaskIntoConstraints="NO" id="976-fk-Kx2">
<rect key="frame" x="172.5" y="597" width="30" height="30"/>
<color key="backgroundColor" white="1" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
<constraints>
......@@ -118,7 +118,7 @@
<constraint firstAttribute="width" constant="30" id="L4p-hP-s5C"/>
</constraints>
</view>
<label opaque="NO" userInteractionEnabled="NO" contentMode="left" horizontalHuggingPriority="251" verticalHuggingPriority="251" text="耗时:" lineBreakMode="tailTruncation" numberOfLines="0" baselineAdjustment="alignBaselines" adjustsFontSizeToFit="NO" translatesAutoresizingMaskIntoConstraints="NO" id="m5L-O7-P31">
<label opaque="NO" userInteractionEnabled="NO" contentMode="left" horizontalHuggingPriority="251" verticalHuggingPriority="251" ambiguous="YES" text="耗时:" lineBreakMode="tailTruncation" numberOfLines="0" baselineAdjustment="alignBaselines" adjustsFontSizeToFit="NO" translatesAutoresizingMaskIntoConstraints="NO" id="m5L-O7-P31">
<rect key="frame" x="15" y="277" width="350" height="38"/>
<constraints>
<constraint firstAttribute="height" constant="38" id="6SS-sb-7I2"/>
......@@ -133,7 +133,7 @@
<constraint firstAttribute="width" secondItem="4ey-Xr-U4e" secondAttribute="height" multiplier="6.5:1" id="8c5-FF-lB9"/>
</constraints>
</imageView>
<textView clipsSubviews="YES" multipleTouchEnabled="YES" contentMode="scaleToFill" editable="NO" text="结果:" textAlignment="natural" translatesAutoresizingMaskIntoConstraints="NO" id="VQn-bS-fWp">
<textView clipsSubviews="YES" multipleTouchEnabled="YES" contentMode="scaleToFill" ambiguous="YES" editable="NO" text="结果:" textAlignment="natural" translatesAutoresizingMaskIntoConstraints="NO" id="VQn-bS-fWp">
<rect key="frame" x="10" y="323" width="355" height="70"/>
<color key="backgroundColor" white="1" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
<constraints>
......@@ -203,6 +203,7 @@
</scene>
</scenes>
<resources>
<image name="hand.jpg" width="564" height="664"/>
<image name="paddle-mobile.png" width="402" height="62"/>
</resources>
</document>
metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift
浏览文件 @ a789f668
......@@ -30,6 +30,7 @@ protocol Net {
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 {
......@@ -39,10 +40,13 @@ extension Net {
getTexture(resTexture)
}
}
func fetchResult(paddleMobileRes: ResultHolder<Float32>) -> [Float32] {
return paddleMobileRes.resultArr
}
}
struct MobileNet: Net{
class MobilenetPreProccess: CusomKernel {
init(device: MTLDevice) {
let s = CusomKernel.Shape.init(inWidth: 224, inHeight: 224, inChannel: 3)
......@@ -100,7 +104,8 @@ struct MobileNet_ssd_hand: Net{
}
func resultStr(res: [Float]) -> String {
fatalError()
return "哈哈哈, 还没好"
// fatalError()
}
func bboxArea(box: [Float32], normalized: Bool) -> Float32 {
......@@ -117,7 +122,6 @@ struct MobileNet_ssd_hand: Net{
}
}
func jaccardOverLap(box1: [Float32], box2: [Float32], normalized: Bool) -> Float32 {
if box2[0] > box1[2] || box2[2] < box1[0] || box2[1] > box1[3] ||
box2[3] < box1[1] {
......@@ -136,9 +140,11 @@ struct MobileNet_ssd_hand: Net{
}
}
func fetchResult(paddleMobileRes: [String : Texture<Float32>]) -> [Float32]{
let bbox = paddleMobileRes["box_coder_0.tmp_0"] ?! " no bbox "
let scores = paddleMobileRes["transpose_12.tmp_0"] ?! " no scores "
func fetchResult(paddleMobileRes: ResultHolder<Float32>) -> [Float32]{
let scores = paddleMobileRes.intermediateResults![0] as! Texture<Float32>
let bbox = paddleMobileRes.intermediateResults![1] as! Texture<Float32>
// let bbox = paddleMobileRes["box_coder_0.tmp_0"] ?! " no bbox "
// let scores = paddleMobileRes["transpose_12.tmp_0"] ?! " no scores "
let score_thredshold: Float32 = 0.01
let nms_top_k = 400
let keep_top_k = 200
......@@ -156,35 +162,44 @@ struct MobileNet_ssd_hand: Net{
var scoreFormatArr: [Float32] = []
var outputArr: [Float32] = []
let numOfOneC = (scores.originDim[2] + 3) / 4 // 480
let cNumOfOneClass = numOfOneC * 4 // 1920
let numOfOneC = (scores.tensorDim[2] + 3) / 4 // 480
let cNumOfOneClass = scores.tensorDim[2] // 1917
let boxSize = bbox.originDim[2] // 4
let classNum = scores.originDim[1] // 7
let cPaddedNumOfOneClass = numOfOneC * 4 // 1920
let boxSize = bbox.tensorDim[2] // 4
let classNum = scores.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 numDet: Int = 0
for i in 0..<classNum {
var sliceScore = scoreFormatArr[(i * cNumOfOneClass)..<((i + 1) * cNumOfOneClass)]
var sliceScore = Array<Float32>(scoreFormatArr[(i * cNumOfOneClass)..<((i + 1) * cNumOfOneClass)])
var scoreThresholdArr: [(Float32, Int)] = []
for i in 0..<cNumOfOneClass {
if sliceScore[i] > score_thredshold {
scoreThresholdArr.append((sliceScore[i], i))
for j in 0..<cNumOfOneClass {
if sliceScore[j] > score_thredshold {
scoreThresholdArr.append((sliceScore[j], j))
}
}
......@@ -204,7 +219,7 @@ struct MobileNet_ssd_hand: Net{
if keep {
let keptIdx = selectedIndex[j].0
let box1 = Array<Float32>(bboxArr[(idx * boxSize)..<(idx * boxSize + 4)])
let box2 = Array<Float32>(bboxArr[(idx * boxSize)..<(keptIdx * boxSize + 4)])
let box2 = Array<Float32>(bboxArr[(keptIdx * boxSize)..<(keptIdx * boxSize + 4)])
let overlap = jaccardOverLap(box1: box1, box2: box2, normalized: true)
keep = (overlap <= nms_threshold)
......@@ -259,7 +274,8 @@ struct MobileNet_ssd_hand: Net{
outputArr.append(contentsOf: subBox)
}
}
print(" fuck success !")
print(outputArr)
return outputArr
}
......
metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift
浏览文件 @ a789f668
......@@ -75,7 +75,7 @@ class ViewController: UIViewController {
}
do {
let max = 10
let max = 1
var startDate = Date.init()
for i in 0..<max {
try inExecutor.predict(input: inTexture, expect: modelHelper.dim, completionHandle: { [weak self] (result) in
......@@ -87,14 +87,16 @@ class ViewController: UIViewController {
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: result.resultArr)
sSelf.resultTextView.text = sSelf.modelHelper.resultStr(res: resultArr)
sSelf.elapsedTimeLabel.text = "平均耗时: \(time/Double(max/2) * 1000.0) ms"
}
}
}, preProcessKernle: self.modelHelper.preprocessKernel)
}, preProcessKernle: self.modelHelper.preprocessKernel, except: 2)
}
} catch let error {
print(error)
......@@ -108,7 +110,7 @@ class ViewController: UIViewController {
threadPickerView.delegate = self
threadPickerView.dataSource = self
selectImage = UIImage.init(named: "banana.jpeg")
selectImage = UIImage.init(named: "hand.jpg")
selectImageView.image = selectImage
modelHelper.getTexture(image: selectImage!.cgImage!) {[weak self] (texture) in
self?.toPredictTexture = texture
......
metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj
浏览文件 @ a789f668
......@@ -46,6 +46,8 @@
FC9D038020E22FBB000F735A /* FeedOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037F20E22FBB000F735A /* FeedOp.swift */; };
FC9D038220E2312E000F735A /* FetchOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D038120E2312E000F735A /* FetchOp.swift */; };
FC9D038420E23B01000F735A /* Texture.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D038320E23B01000F735A /* Texture.swift */; };
FCA3A1632132A4AC00084FE5 /* ReshapeKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCA3A1622132A4AC00084FE5 /* ReshapeKernel.metal */; };
FCA3A1652132A5EB00084FE5 /* Common.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCA3A1642132A5EB00084FE5 /* Common.metal */; };
FCBCCC572122F41300D94F7E /* DwConvBNReluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC562122F41300D94F7E /* DwConvBNReluOp.swift */; };
FCBCCC592122F42700D94F7E /* ConvBNReluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC582122F42700D94F7E /* ConvBNReluOp.swift */; };
FCBCCC5B2122F66F00D94F7E /* ConvBNReluKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC5A2122F66F00D94F7E /* ConvBNReluKernel.swift */; };
......@@ -126,6 +128,8 @@
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>"; };
FC9D038320E23B01000F735A /* Texture.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Texture.swift; sourceTree = "<group>"; };
FCA3A1622132A4AC00084FE5 /* ReshapeKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ReshapeKernel.metal; sourceTree = "<group>"; };
FCA3A1642132A5EB00084FE5 /* Common.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = Common.metal; sourceTree = "<group>"; };
FCBCCC562122F41300D94F7E /* DwConvBNReluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = DwConvBNReluOp.swift; sourceTree = "<group>"; };
FCBCCC582122F42700D94F7E /* ConvBNReluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvBNReluOp.swift; sourceTree = "<group>"; };
FCBCCC5A2122F66F00D94F7E /* ConvBNReluKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvBNReluKernel.swift; sourceTree = "<group>"; };
......@@ -349,6 +353,8 @@
FCDDC6C9212FDF6800E5EF74 /* BatchNormKernel.metal */,
FCDDC6CB212FDFDB00E5EF74 /* ReluKernel.metal */,
FCDDC6CE212FE14700E5EF74 /* PriorBoxKernel.metal */,
FCA3A1622132A4AC00084FE5 /* ReshapeKernel.metal */,
FCA3A1642132A5EB00084FE5 /* Common.metal */,
);
path = metal;
sourceTree = "<group>";
......@@ -482,6 +488,7 @@
FC039BB820E11CC20081E9F8 /* framework.pb.swift in Sources */,
FC039B9920E11C9A0081E9F8 /* Types.swift in Sources */,
FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */,
FCA3A1632132A4AC00084FE5 /* ReshapeKernel.metal in Sources */,
FCBCCC592122F42700D94F7E /* ConvBNReluOp.swift in Sources */,
FC039BA920E11CBC0081E9F8 /* ConvOp.swift in Sources */,
FC9D038420E23B01000F735A /* Texture.swift in Sources */,
......@@ -503,6 +510,7 @@
FCBCCC69212306D300D94F7E /* ConcatKernel.swift in Sources */,
FCDDC6C8212FA3CA00E5EF74 /* ConvTransposeKernel.swift in Sources */,
FC82735920E3C04200BE430A /* OpCreator.swift in Sources */,
FCA3A1652132A5EB00084FE5 /* Common.metal in Sources */,
FCBCCC5D2122F8A100D94F7E /* DepthwiseConvOp.swift in Sources */,
FC0E2DBE20EE460D009C1FAC /* BatchNormKernel.swift in Sources */,
FC039BAB20E11CBC0081E9F8 /* Operator.swift in Sources */,
......
metal/paddle-mobile/paddle-mobile/Common/Extensions.swift
浏览文件 @ a789f668
......@@ -97,6 +97,21 @@ extension Array where Element: Comparable{
}
}
extension Array {
func strideArray(inCount: Int = 20) -> Array<Element> {
if count < inCount {
return self
} else {
let stride = count / inCount
var newArray: [Element] = []
for i in 0..<inCount {
newArray.append(self[i * stride])
}
return newArray
}
}
}
extension String{
func cStr() -> UnsafePointer<Int8>? {
return (self as NSString).utf8String
......
metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift
浏览文件 @ a789f668
......@@ -288,8 +288,8 @@ public extension MTLTexture {
func logDesc<T>(header: String = "", stridable: Bool = true) -> T? {
print(header)
print("texture: \(self)")
// let res: [(index: Int, value: T)] = stridableFloatArray(stridable: stridable)
// print(res)
// let res: [(index: Int, value: T)] = stridableFloatArray(stridable: stridable)
// print(res)
if textureType == .type2DArray {
for i in 0..<arrayLength{
......@@ -301,8 +301,8 @@ public extension MTLTexture {
getBytes(bytes, bytesPerRow: bytesPerRow, bytesPerImage: bytesPerImage, from: region, mipmapLevel: 0, slice: i)
let p = bytes.assumingMemoryBound(to: T.self)
str += "2d array count : \(width * height * depth * 4) \n"
if stridable && width * height * depth * 4 > 100 {
for j in stride(from: 0, to: width * height * depth * 4 , by: width * height * depth * 4 / 100){
if stridable && width * height * depth * 4 > 20 {
for j in stride(from: 0, to: width * height * depth * 4 , by: width * height * depth * 4 / 20){
str += " index \(j): \(p[j])"
}
} else {
......@@ -324,7 +324,7 @@ public extension MTLTexture {
str += "2d count : \(width * width * 4) \n"
if stridable {
for j in stride(from: 0, to: width * height * 4, by: width * height * 4 / 100){
for j in stride(from: 0, to: width * height * 4, by: width * height * 4 / 20){
str += "index \(j): \(p[j]) "
}
} else {
......@@ -375,7 +375,14 @@ public extension MTLBuffer {
return texture
}
func array<T>() -> [T] {
var array: [T] = []
let pointer = contents().bindMemory(to: T.self, capacity: length)
for i in 0..<(length / MemoryLayout<T>.size) {
array.append(pointer[i])
}
return array;
}
}
......
metal/paddle-mobile/paddle-mobile/Common/Types.swift
浏览文件 @ a789f668
......@@ -202,7 +202,10 @@ extension DataLayout: Equatable {
if lhs.layoutWithDim.count == rhs.layoutWithDim.count {
var result = true
for i in 0..<lhs.layoutWithDim.count {
result = (lhs.layoutWithDim[i] == rhs.layoutWithDim[i])
result = (lhs.layoutWithDim[i].0 == rhs.layoutWithDim[i].0)
if !result {
break
}
}
return result
} else {
......@@ -215,7 +218,7 @@ extension DataLayout: Equatable {
protocol Variant: CustomStringConvertible, CustomDebugStringConvertible {
public protocol Variant: CustomStringConvertible, CustomDebugStringConvertible {
}
extension Tensor: Variant {
......
metal/paddle-mobile/paddle-mobile/Executor.swift
浏览文件 @ a789f668
......@@ -17,9 +17,9 @@ import Foundation
public class ResultHolder<P: PrecisionType> {
public let dim: [Int]
public let resultArr: [P]
public var intermediateResults: [Texture<P>]?
public var intermediateResults: [Variant]?
public let elapsedTime: Double
public init(inDim: [Int], inResult: [P], inElapsedTime: Double, inIntermediateResults: [Texture<P>]? = nil) {
public init(inDim: [Int], inResult: [P], inElapsedTime: Double, inIntermediateResults: [Variant]? = nil) {
dim = inDim
resultArr = inResult
elapsedTime = inElapsedTime
......@@ -60,11 +60,11 @@ public class Executor<P: PrecisionType> {
queue = inQueue
for block in inProgram.programDesc.blocks {
//block.ops.count
for i in 0..<block.ops.count {
for i in 0..<39 {
let op = block.ops[i]
do {
let op = try OpCreator<P>.shared.creat(device: inDevice, opDesc: op, scope: inProgram.scope)
op.inferShape()
// op.inferShape()
ops.append(op)
} catch let error {
throw error
......@@ -73,7 +73,7 @@ public class Executor<P: PrecisionType> {
}
}
public func predict(input: MTLTexture, expect: [Int], completionHandle: @escaping (ResultHolder<P>) -> Void, preProcessKernle: CusomKernel? = nil) throws {
public func predict(input: MTLTexture, expect: [Int], completionHandle: @escaping (ResultHolder<P>) -> Void, preProcessKernle: CusomKernel? = nil, except: Int = 0) throws {
guard let buffer = queue.makeCommandBuffer() else {
throw PaddleMobileError.predictError(message: "CommandBuffer is nil")
}
......@@ -92,8 +92,9 @@ public class Executor<P: PrecisionType> {
let beforeDate = Date.init()
let inputTexture = InputTexture.init(inMTLTexture: resInput, inExpectDim: Dim.init(inDim: expect))
program.scope.setInput(input: inputTexture)
for op in ops {
//(ops.count - except)
for i in 0..<ops.count {
let op = ops[i]
do {
try op.run(device: device, buffer: buffer)
} catch let error {
......@@ -101,36 +102,55 @@ public class Executor<P: PrecisionType> {
}
}
var outputTextures: [Variant]?
if except > 0 {
outputTextures = ops[ops.count - except].inputs()
}
buffer.addCompletedHandler { (commandbuffer) in
// return;
// let inputArr = resInput.floatArray(res: { (p:P) -> P in
// return p
// })
// print(inputArr)
// let stridableInput: [(index: Int, value: Float)] = input.stridableFloatArray()
// print(stridableInput)
// writeToLibrary(fileName: "input_hand", array: inputArr)
// 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 op in self.ops {
// let _: Flo? = input.logDesc(header: "input: ", stridable: true)
for op in self.ops {
// op.delogOutput()
// }
}
// return
// self.ops[2].delogOutput()
self.ops[38].delogOutput()
// self.ops[91].delogOutput()
// self.ops[92].delogOutput()
// self.ops[93].delogOutput()
return;
let afterDate = Date.init()
guard let outputVar = self.program.scope.output() else {
fatalError("output nil")
}
var resultHolder: ResultHolder<P>
if except > 0 {
resultHolder = ResultHolder<P>.init(inDim: [], inResult: [], inElapsedTime: afterDate.timeIntervalSince(beforeDate), inIntermediateResults: outputTextures)
} else {
let outputVar: Variant = self.program.scope.output()!
let output: Texture<P> = outputVar as! Texture<P>
guard let output = outputVar as? Texture<P> else {
fatalError("output var type error")
}
let resultHodlder = ResultHolder<P>.init(inDim: output.dim.dims, inResult: output.metalTexture.floatArray(res: { (p:P) -> P in
resultHolder = ResultHolder<P>.init(inDim: output.dim.dims, inResult: output.metalTexture.floatArray(res: { (p:P) -> P in
return p
}), inElapsedTime: afterDate.timeIntervalSince(beforeDate))
completionHandle(resultHodlder)
}
completionHandle(resultHolder)
}
buffer.commit()
}
......
metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift
浏览文件 @ a789f668
......@@ -25,6 +25,7 @@ protocol Runable {
func run(device: MTLDevice, buffer: MTLCommandBuffer) throws
func runImpl(device: MTLDevice,buffer: MTLCommandBuffer) throws
func delogOutput()
func inputs() -> [Variant]
}
extension Runable where Self: OperatorProtocol{
......
metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift
浏览文件 @ a789f668
......@@ -43,6 +43,11 @@ class BatchNormParam<P: PrecisionType>: OpParam {
}
class BatchNormOp<P: PrecisionType>: Operator<BatchNormKernel<P>, BatchNormParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input, para.inputBias, para.inputMean, para.inputScale, para.inputVariance]
}
func inferShape() {
para.output.dim = para.input.dim
}
......
metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift
浏览文件 @ a789f668
///* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License. */
import Foundation
class BoxcoderParam<P: PrecisionType>: OpParam {
///* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License. */
import Foundation
class BoxcoderParam<P: PrecisionType>: OpParam {
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
......@@ -39,12 +39,16 @@ class BoxcoderParam<P: PrecisionType>: OpParam {
var output: Texture<P>
let codeType: String
let boxNormalized: Bool
}
}
class BoxcoderOp<P: PrecisionType>: Operator<BoxcoderKernel<P>, BoxcoderParam<P>>, Runable, Creator, InferShaperable{
class BoxcoderOp<P: PrecisionType>: Operator<BoxcoderKernel<P>, BoxcoderParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.priorBox, para.priorBoxVar, para.targetBox]
}
func inferShape() {
// para.output.dim = para.input.dim
// para.output.dim = para.input.dim
}
typealias OpType = BoxcoderOp<P>
......@@ -55,7 +59,22 @@ class BoxcoderOp<P: PrecisionType>: Operator<BoxcoderKernel<P>, BoxcoderParam<P>
throw error
}
}
}
func delogOutput() {
let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in
return o
}
print(outputArray.strideArray())
//box_coder_0.tmp_0
// writeToLibrary(fileName: "boxcoder_output", array: outputArray)
print(para.output.metalTexture)
print(" write done ")
}
}
metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift
浏览文件 @ a789f668
......@@ -40,9 +40,13 @@ class ConcatParam<P: PrecisionType>: OpParam {
class ConcatOp<P: PrecisionType>: Operator<ConcatKernel<P>, ConcatParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return para.input
}
func inferShape() {
// let dim = para.input.reduce([0, 0]) {[$0[0] + $1.dim[0], $1.dim[1]]}
// para.output.dim = Dim.init(inDim: dim)
// let dim = para.input.reduce([0, 0]) {[$0[0] + $1.dim[0], $1.dim[1]]}
// para.output.dim = Dim.init(inDim: dim)
}
typealias OpType = ConcatOp<P>
......@@ -53,6 +57,25 @@ class ConcatOp<P: PrecisionType>: Operator<ConcatKernel<P>, ConcatParam<P>>, Run
throw error
}
}
func delogOutput() {
let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in
return o
}
print(outputArray.strideArray())
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())
}
}
......
metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift
浏览文件 @ a789f668
......@@ -60,6 +60,11 @@ class ConvAddBatchNormReluParam<P: PrecisionType>: OpParam {
}
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>
func inferShape() {
......@@ -111,25 +116,25 @@ class ConvAddBatchNormReluOp<P: PrecisionType>: Operator<ConvAddBatchNormReluKer
func delogOutput() {
// let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false)
// para.filter.logDataPointer(header: "filter data pointer: ")
// print("filter: \(para.filter)")
// let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false)
// para.filter.logDataPointer(header: "filter data pointer: ")
// print("filter: \(para.filter)")
// print("biase: \(para.y)")
// print("padding: \(para.paddings)")
// print("stride: \(para.stride)")
// 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 _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false)
// let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false)
// let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false)
let output = para.output.metalTexture.floatArray { (p: P) -> P in
return p
}
//
//
writeToLibrary(fileName: "output_112x112x32_2", array: output)
print(" write done")
// let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false)
// let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false)
}
}
metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift
浏览文件 @ a789f668
......@@ -43,6 +43,29 @@ class ConvAddParam<P: PrecisionType>: OpParam {
}
class ConvAddOp<P: PrecisionType>: Operator<ConvAddKernel<P>, ConvAddParam<P>>, Runable, Creator, InferShaperable, Fusion{
func delogOutput() {
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 {
let beginNode = Node.init(inType: gConvType)
_ = beginNode
......@@ -54,6 +77,11 @@ class ConvAddOp<P: PrecisionType>: Operator<ConvAddKernel<P>, ConvAddParam<P>>,
return [:]
}
func inputs() -> [Variant] {
return [para.input, para.y, para.filter]
}
static func fusionType() -> String {
return gConvAddType
}
......
metal/paddle-mobile/paddle-mobile/Operators/ConvBNReluOp.swift
浏览文件 @ a789f668
......@@ -57,6 +57,11 @@ class ConvBNReluParam<P: PrecisionType>: OpParam {
class ConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNReluParam<P>>, Runable, Creator, InferShaperable, Fusion{
typealias OpType = ConvBNReluOp<P>
func inputs() -> [Variant] {
return [para.input, para.variance, para.bias, para.mean, para.scale, para.filter]
}
func inferShape() {
let inDims = para.input.dim
let filterDim = para.filter.dim
......@@ -117,13 +122,57 @@ class ConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNReluPa
// 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
}
//
writeToLibrary(fileName: "output_112x112x32_2", array: output)
print(output)
//
writeToLibrary(fileName: "batch_norm_34.tmp_2", array: output)
print(" write done")
//
// let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false)
// let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: true)
}
}
metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift
浏览文件 @ a789f668
......@@ -41,6 +41,11 @@ class ConvParam<P: PrecisionType>: OpParam {
}
class ConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable, Creator, InferShaperable {
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)
......@@ -83,6 +88,6 @@ class ConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable,
func delogOutput() {
print("conv output : ")
print(para.output.metalTexture)
// let _: Float16? = para.output.metalTexture.logDesc()
// let _: Float16? = para.output.metalTexture.logDesc()
}
}
metal/paddle-mobile/paddle-mobile/Operators/ConvTransposeOp.swift
浏览文件 @ a789f668
......@@ -28,6 +28,10 @@ class ConvTransposeParam<P: PrecisionType>: ConvParam<P> {
class ConvTransposeOp<P: PrecisionType>: Operator<ConvTransposeKernel<P>, ConvTransposeParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input, para.filter]
}
func inferShape() {
// para.output.dim = para.input.dim
}
......
metal/paddle-mobile/paddle-mobile/Operators/DepthwiseConvOp.swift
浏览文件 @ a789f668
......@@ -15,6 +15,11 @@
import Foundation
class DepthConvOp<P: PrecisionType>: Operator<ConvKernel<P>, ConvParam<P>>, Runable, Creator, InferShaperable {
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)
......
metal/paddle-mobile/paddle-mobile/Operators/DwConvBNReluOp.swift
浏览文件 @ a789f668
......@@ -17,6 +17,10 @@ import Foundation
class DwConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNReluParam<P>>, Runable, Creator, InferShaperable, Fusion{
typealias OpType = ConvBNReluOp<P>
func inputs() -> [Variant] {
return [para.input, para.bias, para.mean, para.filter, para.variance, para.scale]
}
func inferShape() {
let inDims = para.input.dim
let filterDim = para.filter.dim
......@@ -77,12 +81,12 @@ class DwConvBNReluOp<P: PrecisionType>: Operator<ConvBNReluKernel<P>, ConvBNRelu
// let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false)
// let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false)
let output = para.output.metalTexture.floatArray { (p: P) -> P in
return p
}
//
writeToLibrary(fileName: "output_112x112x32_2", array: output)
print(" write done")
// let 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)
}
......
metal/paddle-mobile/paddle-mobile/Operators/ElementwiseAddOp.swift
浏览文件 @ a789f668
......@@ -35,6 +35,10 @@ class ElementwiseAddParam<P: PrecisionType>: OpParam {
class ElementwiseAddOp<P: PrecisionType>: Operator<ElementwiseAddKernel<P>, ElementwiseAddParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input, para.inputY]
}
func inferShape() {
para.output.dim = para.input.dim
}
......
metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift
浏览文件 @ a789f668
......@@ -36,6 +36,10 @@ class FeedParam<P: PrecisionType>: OpParam{
class FeedOp<P: PrecisionType>: Operator<Texture2DTo2DArrayKernel<P>, FeedParam<P>>, Runable, Creator, InferShaperable {
typealias OpType = FeedOp<P>
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() {
// print("feed input: \(para.input.expectDim)")
print("feed output: \(para.output.dim)")
......@@ -50,19 +54,19 @@ class FeedOp<P: PrecisionType>: Operator<Texture2DTo2DArrayKernel<P>, FeedParam<
throw error
}
// let resizeKernel = ResizeKernel<P>.init(device: device)
// let resizeParam = ResizeParam.init(input: para.input.mtlTexture, output: para.output.metalTexture, expectDim: para.input.expectDim)
// do {
// try resizeKernel.compute(commandBuffer: buffer, param: resizeParam)
// } catch let error {
// throw error
// }
// let resizeKernel = ResizeKernel<P>.init(device: device)
// let resizeParam = ResizeParam.init(input: para.input.mtlTexture, output: para.output.metalTexture, expectDim: para.input.expectDim)
// do {
// try resizeKernel.compute(commandBuffer: buffer, param: resizeParam)
// } catch let error {
// throw error
// }
}
func delogOutput() {
// para.input.mtlTexture.logDesc()
// let _: P? = para.input.mtlTexture.logDesc(header: "feed input: ", stridable: true)
// let _: P? = para.output.metalTexture.logDesc(header: "feed output: ", stridable: false)
// para.input.mtlTexture.logDesc()
// let _: P? = para.input.mtlTexture.logDesc(header: "feed input: ", stridable: true)
// let _: P? = para.output.metalTexture.logDesc(header: "feed output: ", stridable: false)
}
}
metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift
浏览文件 @ a789f668
......@@ -42,6 +42,10 @@ class FetchKernel<P: PrecisionType>: Kernel, Computable {
}
class FetchOp<P: PrecisionType>: Operator< FetchKernel<P>, FetchParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() {
print(para.input.dim)
}
......
metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift
浏览文件 @ a789f668
......@@ -50,7 +50,7 @@ class ConvAddBatchNormReluKernel<P: PrecisionType>: Kernel, Computable, Testable
required init(device: MTLDevice, param: ConvAddBatchNormReluParam<P>) {
param.output.initTexture(device: device, transpose: [0, 2, 3, 1])
param.output.initTexture(device: device, inTranspose: [0, 2, 3, 1])
if param.filter.width == 1 && param.filter.height == 1 {
super.init(device: device, inFunctionName: "conv_add_batch_norm_relu_1x1")
......
metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift
浏览文件 @ a789f668
......@@ -25,7 +25,7 @@ class ConvAddKernel<P: PrecisionType>: Kernel, Computable {
super.init(device: device, inFunctionName: "conv_add_3x3")
}
param.output.initTexture(device: device, transpose: [0, 3, 1, 2])
param.output.initTexture(device: device, inTranspose: [0, 3, 2, 1])
let offsetX = param.filter.width/2 - Int(param.paddings[0])
let offsetY = param.filter.height/2 - Int(param.paddings[1])
......
metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift
浏览文件 @ a789f668
......@@ -59,7 +59,7 @@ class ConvBNReluKernel<P: PrecisionType>: Kernel, Computable, Testable {
} else {
super.init(device: device, inFunctionName: "conv_batch_norm_relu_3x3")
}
param.output.initTexture(device: device, transpose: [0, 2, 3, 1])
param.output.initTexture(device: device, inTranspose: [0, 2, 3, 1])
param.filter.initBuffer(device: device, precision: Tensor.BufferPrecision.Float32)
param.variance.initBuffer(device: device)
......@@ -70,8 +70,13 @@ class ConvBNReluKernel<P: PrecisionType>: Kernel, Computable, Testable {
let offsetX = param.filter.width/2 - Int(param.paddings[0])
let offsetY = param.filter.height/2 - Int(param.paddings[1])
print("offset x: \(offsetX)")
print("offset y: \(offsetY)")
print(" param filter width: \(param.filter.width)")
print(" param filter height: \(param.filter.height)")
print(" param paddings: \(param.paddings)")
print("ConvBNReluKernel offset x: \(offsetX)")
print("ConvBNReluKernel offset y: \(offsetY)")
let offsetZ = 0.0
......@@ -116,8 +121,8 @@ class ConvBNReluKernel<P: PrecisionType>: Kernel, Computable, Testable {
encoder.setTexture(param.output.metalTexture, index: 1)
encoder.setBytes(&metalParam, length: MemoryLayout<MetalConvParam>.size, index: 0)
encoder.setBuffer(param.filter.buffer, offset: 0, index: 1)
encoder.setBuffer(param.newScale!, offset: 0, index: 3)
encoder.setBuffer(param.newBiase!, offset: 0, index: 4)
encoder.setBuffer(param.newScale!, offset: 0, index: 2)
encoder.setBuffer(param.newBiase!, offset: 0, index: 3)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
}
......@@ -132,9 +137,8 @@ class ConvBNReluKernel<P: PrecisionType>: Kernel, Computable, Testable {
var inMetalParam = param.metalParam
encoder.setBytes(&inMetalParam, length: MemoryLayout<MetalConvParam>.size, index: 0)
encoder.setBuffer(param.filterBuffer, offset: 0, index: 1)
encoder.setBuffer(param.biaseBuffer, offset: 0, index: 2)
encoder.setBuffer(param.newScaleBuffer, offset: 0, index: 3)
encoder.setBuffer(param.newBiaseBuffer, offset: 0, index: 4)
encoder.setBuffer(param.newScaleBuffer, offset: 0, index: 2)
encoder.setBuffer(param.newBiaseBuffer, offset: 0, index: 3)
encoder.dispatch(computePipline: pipline, outTexture: param.outputTexture)
encoder.endEncoding()
}
......
metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift
浏览文件 @ a789f668
......@@ -34,14 +34,14 @@ class PriorBoxKernel<P: PrecisionType>: Kernel, Computable{
required init(device: MTLDevice, param: PriorBoxParam<P>) {
super.init(device: device, inFunctionName: "prior_box")
param.output.initTexture(device: device, transpose: [2, 0, 1, 3])
param.outputVariances.initTexture(device: device, transpose: [2, 0, 1, 3])
param.output.initTexture(device: device, inTranspose: [2, 0, 1, 3])
param.outputVariances.initTexture(device: device, inTranspose: [2, 0, 1, 3])
let imageWidth = Float32(param.inputImage.originDim[3])
let imageHeight = Float32(param.inputImage.originDim[2])
let featureWidth = param.inputImage.originDim[3]
let featureHeight = param.inputImage.originDim[2]
let featureWidth = param.input.originDim[3]
let featureHeight = param.input.originDim[2]
if param.stepW == 0 || param.stepH == 0 {
param.stepW = Float32(imageWidth) / Float32(featureWidth)
......@@ -88,11 +88,22 @@ class PriorBoxKernel<P: PrecisionType>: Kernel, Computable{
guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
throw PaddleMobileError.predictError(message: " encode is nil")
}
print("metalParam: \(metalParam)")
print(" newAspectRatios ")
print(param.newAspectRatios!)
print(" clip: \(metalParam.clip)")
print(" metalParam.numPriors: \(metalParam.numPriors)")
print(" aspecRatiosSize: \(metalParam.aspecRatiosSize)")
encoder.setTexture(param.input.metalTexture, index: 0)
encoder.setTexture(param.output.metalTexture, index: 1)
encoder.setTexture(param.outputVariances.metalTexture, index: 2)
encoder.setBytes(&metalParam, length: MemoryLayout<PriorBoxMetalParam>.size, index: 0)
encoder.setBytes(param.aspectRatios, length: MemoryLayout<Float32>.size * param.aspectRatios.count, index: 1)
encoder.setBytes(param.newAspectRatios!, length: MemoryLayout<Float32>.size * param.newAspectRatios!.count, index: 1)
encoder.setBytes(param.variances, length: MemoryLayout<Float32>.size * param.variances.count, index: 2)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
......
metal/paddle-mobile/paddle-mobile/Operators/Kernels/Texture2DTo2DArrayKernel.swift
浏览文件 @ a789f668
......@@ -32,7 +32,7 @@ class Texture2DTo2DArrayKernel<P: PrecisionType>: Kernel, Computable{
}
required init(device: MTLDevice, param: FeedParam<P>) {
param.output.initTexture(device: device, transpose: [0, 2, 3, 1])
param.output.initTexture(device: device, inTranspose: [0, 2, 3, 1])
super.init(device: device, inFunctionName: "texture2d_to_2d_array")
}
}
metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift
浏览文件 @ a789f668
......@@ -41,15 +41,35 @@ struct TransposeTestParam: TestParam {
}
class TransposeKernel<P: PrecisionType>: Kernel, Computable, Testable {
var metalParam: TransposeMetalParam!
func compute(commandBuffer: MTLCommandBuffer, param: TransposeParam<P>) throws {
guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
throw PaddleMobileError.predictError(message: " encode is nil")
}
encoder.setTexture(param.input.metalTexture, index: 0)
encoder.setTexture(param.output.metalTexture, index: 1)
encoder.setBytes(&metalParam, length: MemoryLayout<TransposeMetalParam>.size, index: 0)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
}
required init(device: MTLDevice, param: TransposeParam<P>) {
param.output.initTexture(device: device, inTranspose: [0, 1, 2, 3])
super.init(device: device, inFunctionName: "transpose")
var invT: [Int] = [0, 1, 2, 3]
for (i, v) in param.input.transpose.enumerated() {
invT[v] = i
}
var axis: [Int] = [0, 1, 2, 3]
// var doNothing = false
// if param.axis.count == param.input.transpose.count {
// doNothing = param.axis == param.input.transpose.map { Int32($0) }
// }
for i in 0..<param.axis.count {
axis[4-param.axis.count+i] = 4 - param.axis.count + Int(param.axis[i])
}
......@@ -62,19 +82,11 @@ class TransposeKernel<P: PrecisionType>: Kernel, Computable, Testable {
} else {
print("====> transpose! SLOW :(")
}
encoder.setTexture(param.input.metalTexture, index: 0)
encoder.setTexture(param.output.metalTexture, index: 1)
encoder.setBytes(&tmp, length: MemoryLayout<TransposeMetalParam>.size, index: 0)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
}
required init(device: MTLDevice, param: TransposeParam<P>) {
param.output.initTexture(device: device, transpose: [0, 1, 2, 3])
super.init(device: device, inFunctionName: "transpose")
metalParam = tmp
}
required init(device: MTLDevice, testParam: TransposeTestParam) {
super.init(device: device, inFunctionName: "transpose")
fatalError()
}
public func test(commandBuffer: MTLCommandBuffer, param: TransposeTestParam) {
......
metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Common.metal 0 → 100644
浏览文件 @ a789f668
//
// common.metal
// paddle-mobile
//
// Created by liuRuiLong on 2018/8/26.
// Copyright © 2018年 orange. All rights reserved.
//
#include <metal_stdlib>
using namespace metal;
inline void xyzn2abcd(int C, int xyzn[4], int abcd[4]) {
abcd[2] = xyzn[0];
abcd[1] = xyzn[1];
uint t = xyzn[2] * 4 + xyzn[3];
abcd[0] = t / C;
abcd[3] = t % C;
}
inline void abcd2xyzn(int C, int abcd[4], int xyzn[4]) {
xyzn[0] = abcd[2];
xyzn[1] = abcd[1];
uint t = abcd[0] * C + abcd[3];
xyzn[2] = t / 4;
xyzn[3] = t % 4;
}
inline int32_t abcd2index(int32_t dim[4], int32_t abcd[4]) {
int32_t r = abcd[0];
r = r * dim[1] + abcd[1];
r = r * dim[2] + abcd[2];
r = r * dim[3] + abcd[3];
return r;
}
inline void index2abcd(int32_t dim[4], int32_t ind, int32_t abcd[4]) {
abcd[3] = ind % dim[3]; ind /= dim[3];
abcd[2] = ind % dim[2]; ind /= dim[2];
abcd[1] = ind % dim[1]; ind /= dim[1];
abcd[0] = ind;
}
inline void trans(int32_t trans[4], int32_t ipos[4], int32_t opos[4]) {
for (int i = 0; i < 4; i++) {
opos[i] = ipos[trans[i]];
}
}
inline void invtrans(int32_t trans[4], int32_t ipos[4], int32_t opos[4]) {
for (int i = 0; i < 4; i++) {
opos[trans[i]] = ipos[i];
}
}
metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ConvKernel.metal
浏览文件 @ a789f668
......@@ -704,9 +704,8 @@ kernel void conv_batch_norm_relu_1x1(texture2d_array<float, access::sample> inTe
texture2d_array<float, access::write> outTexture [[texture(1)]],
constant MetalConvParam &param [[buffer(0)]],
const device float4 *weights [[buffer(1)]],
const device float4 *biase [[buffer(2)]],
const device float4 *new_scale [[buffer(3)]],
const device float4 *new_biase [[buffer(4)]],
const device float4 *new_scale [[buffer(2)]],
const device float4 *new_biase [[buffer(3)]],
uint3 gid [[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() ||
......@@ -749,9 +748,8 @@ kernel void conv_batch_norm_relu_3x3(texture2d_array<float, access::sample> inTe
texture2d_array<float, access::write> outTexture [[texture(1)]],
constant MetalConvParam &param [[buffer(0)]],
const device float4 *weights [[buffer(1)]],
const device float4 *biase [[buffer(2)]],
const device float4 *new_scale [[buffer(3)]],
const device float4 *new_biase [[buffer(4)]],
const device float4 *new_scale [[buffer(2)]],
const device float4 *new_biase [[buffer(3)]],
uint3 gid [[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() ||
......@@ -803,8 +801,8 @@ kernel void depthwise_conv_batch_norm_relu_3x3(texture2d_array<float, access::sa
texture2d_array<float, access::write> outTexture [[texture(1)]],
constant MetalConvParam &param [[buffer(0)]],
const device float *weights [[buffer(1)]],
const device float4 *new_scale [[buffer(3)]],
const device float4 *new_biase [[buffer(4)]],
const device float4 *new_scale [[buffer(2)]],
const device float4 *new_biase [[buffer(3)]],
uint3 gid [[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() ||
......
metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Kernels.metal
浏览文件 @ a789f668
......@@ -13,6 +13,7 @@
limitations under the License. */
#include <metal_stdlib>
#include "Common.metal"
using namespace metal;
struct OutputDim {
......@@ -206,48 +207,7 @@ kernel void softmax_half(texture2d_array<half, access::read> inTexture [[texture
outTexture.write(rr, gid.xy, gid.z);
}
inline void xyzn2abcd(int C, int xyzn[4], int abcd[4]) {
abcd[2] = xyzn[0];
abcd[1] = xyzn[1];
uint t = xyzn[2] * 4 + xyzn[3];
abcd[0] = t / C;
abcd[3] = t % C;
}
inline void abcd2xyzn(int C, int abcd[4], int xyzn[4]) {
xyzn[0] = abcd[2];
xyzn[1] = abcd[1];
uint t = abcd[0] * C + abcd[3];
xyzn[2] = t / 4;
xyzn[3] = t % 4;
}
inline int32_t abcd2index(int32_t dim[4], int32_t abcd[4]) {
int32_t r = abcd[0];
r = r * dim[1] + abcd[1];
r = r * dim[2] + abcd[2];
r = r * dim[3] + abcd[3];
return r;
}
inline void index2abcd(int32_t dim[4], int32_t ind, int32_t abcd[4]) {
abcd[3] = ind % dim[3]; ind /= dim[3];
abcd[2] = ind % dim[2]; ind /= dim[2];
abcd[1] = ind % dim[1]; ind /= dim[1];
abcd[0] = ind;
}
inline void trans(int32_t trans[4], int32_t ipos[4], int32_t opos[4]) {
for (int i = 0; i < 4; i++) {
opos[i] = ipos[trans[i]];
}
}
inline void invtrans(int32_t trans[4], int32_t ipos[4], int32_t opos[4]) {
for (int i = 0; i < 4; i++) {
opos[trans[i]] = ipos[i];
}
}
struct TransposeParam {
int iC;
......@@ -260,6 +220,7 @@ kernel void transpose(texture2d_array<float, access::read> inTexture [[texture(0
constant TransposeParam &pm [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
if ((pm.axis[0] == 0) && (pm.axis[1] == 1) && (pm.axis[2] == 2) && (pm.axis[3] == 3)) {
// do nothing
float4 r = inTexture.read(gid.xy, gid.z);
......@@ -282,56 +243,6 @@ kernel void transpose(texture2d_array<float, access::read> inTexture [[texture(0
}
}
struct ReshapeParam {
int32_t idim[4];
int32_t itrans[4];
int32_t odim[4];
int32_t otrans[4];
};
kernel void reshape(texture2d_array<float, access::read> inTexture [[texture(0)]],
texture2d_array<float, access::write> outTexture [[texture(1)]],
constant ReshapeParam &rp [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() ||
gid.y >= outTexture.get_height() ||
gid.z >= outTexture.get_array_size()) return;
int oxyzn[4] = {int(gid.x), int(gid.y), int(gid.z), 0}, oabcd[4], ixyzn[4], iabcd[4];
ReshapeParam lrp = rp;
int oC = lrp.odim[lrp.otrans[3]];
int iC = lrp.idim[lrp.itrans[3]];
int count = lrp.odim[0] * lrp.odim[1] * lrp.odim[2] * lrp.odim[3];
float4 r;
for (int n = 0; n < 4; n++) {
oxyzn[3] = n;
xyzn2abcd(oC, oxyzn, oabcd);
int tabcd[4];
invtrans(lrp.otrans, oabcd, tabcd);
int index = abcd2index(lrp.odim, tabcd);
if (index < count) {
index2abcd(lrp.idim, index, tabcd);
trans(lrp.itrans, tabcd, iabcd);
abcd2xyzn(iC, tabcd, ixyzn);
r[n] = inTexture.read(uint2(ixyzn[0], ixyzn[1]), ixyzn[2])[ixyzn[3]];
} else {
r[n] = 0;
}
}
outTexture.write(r, gid.xy, gid.z);
}
//
//kernel void reshape_half(texture2d_array<half, access::read> inTexture [[texture(0)]],
// texture2d_array<half, access::write> outTexture [[texture(1)]],
// uint3 gid [[thread_position_in_grid]]) {
// if (gid.x >= outTexture.get_width() ||
// gid.y >= outTexture.get_height() ||
// gid.z >= outTexture.get_array_size()) return;
//
// half4 r = inTexture.read(uint2(0, 0), gid.x);
// outTexture.write(r, gid.xy, gid.z);
//}
struct ConcatParam {
int32_t odim[4];
int32_t axis;
......
metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PriorBoxKernel.metal
浏览文件 @ a789f668
......@@ -60,7 +60,7 @@ kernel void prior_box(texture2d_array<float, access::read> inTexture [[texture(0
float4 res;
if (param.clip) {
res = min(max(box, 0.0), 1.0);
res = fmin(fmax(box, 0.0), 1.0);
} else {
res = box;
}
......@@ -92,6 +92,7 @@ kernel void prior_box(texture2d_array<float, access::read> inTexture [[texture(0
variances_output.y = variance.y;
variances_output.z = variance.z;
variances_output.w = variance.w;
varianceTexture.write(variances_output, gid.xy, gid.z);
}
}
metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal 0 → 100644
浏览文件 @ a789f668
/* 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. */
#include <metal_stdlib>
#include "Common.metal"
using namespace metal;
struct ReshapeParam {
int32_t idim[4];
int32_t itrans[4];
int32_t odim[4];
int32_t otrans[4];
};
kernel void reshape(texture2d_array<float, access::read> inTexture [[texture(0)]],
texture2d_array<float, access::write> outTexture [[texture(1)]],
constant ReshapeParam &rp [[buffer(0)]],
uint3 gid [[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() ||
gid.y >= outTexture.get_height() ||
gid.z >= outTexture.get_array_size()) return;
int oxyzn[4] = {int(gid.x), int(gid.y), int(gid.z), 0}, oabcd[4], ixyzn[4];
ReshapeParam lrp = rp;
int oC = lrp.odim[lrp.otrans[3]];
int iC = lrp.idim[lrp.itrans[3]];
int count = lrp.odim[0] * lrp.odim[1] * lrp.odim[2] * lrp.odim[3];
float4 r;
for (int n = 0; n < 4; n++) {
oxyzn[3] = n;
//4 (gid.x gid.y, gid.z, 0~4)
xyzn2abcd(oC, oxyzn, oabcd);
int tabcd[4];
invtrans(lrp.otrans, oabcd, tabcd);
int index = abcd2index(lrp.odim, tabcd);
if (index < count) {
int c = index % 4;
int temp0 = index % (inTexture.get_array_size() * 4);
int slice = temp0 / 4;
int temp1 = index % (inTexture.get_array_size() * 4 * lrp.idim[2]);
int w = temp1 / (inTexture.get_array_size() * 4);
int h = index / (inTexture.get_array_size() * 4 * lrp.idim[2]);
// index2abcd(lrp.idim, index, tabcd);
// abcd2xyzn(iC, tabcd, ixyzn);
r[n] = inTexture.read(uint2(w, h), slice)[c];
} else {
r[n] = 0;
}
}
outTexture.write(r, gid.xy, gid.z);
}
//
//kernel void reshape_half(texture2d_array<half, access::read> inTexture [[texture(0)]],
// texture2d_array<half, access::write> outTexture [[texture(1)]],
// uint3 gid [[thread_position_in_grid]]) {
// if (gid.x >= outTexture.get_width() ||
// gid.y >= outTexture.get_height() ||
// gid.z >= outTexture.get_array_size()) return;
//
// half4 r = inTexture.read(uint2(0, 0), gid.x);
// outTexture.write(r, gid.xy, gid.z);
//}
metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift
浏览文件 @ a789f668
......@@ -32,6 +32,10 @@ class MulticlassNMSParam<P: PrecisionType>: OpParam {
class MulticlassNMSOp<P: PrecisionType>: Operator<MulticlassNMSKernel<P>, MulticlassNMSParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.scores,para.bboxes]
}
func inferShape() {
// para.output.dim = para.input.dim
}
......
metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift
浏览文件 @ a789f668
......@@ -29,7 +29,7 @@ class PoolParam<P: PrecisionType>: OpParam {
} catch let error {
throw error
}
// let buffer = input.metalTexture.buffer.contents().assumingMemoryBound(to: P.self)
// let buffer = input.metalTexture.buffer.contents().assumingMemoryBound(to: P.self)
}
let input: Texture<P>
var output: Texture<P>
......@@ -43,6 +43,10 @@ class PoolParam<P: PrecisionType>: OpParam {
class PoolOp<P: PrecisionType>: Operator<PoolKernel<P>, PoolParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() {
// para.output.dim = para.input.dim
}
......
metal/paddle-mobile/paddle-mobile/Operators/PreluOp.swift
浏览文件 @ a789f668
......@@ -35,6 +35,10 @@ class PreluParam<P: PrecisionType>: OpParam {
class PreluOp<P: PrecisionType>: Operator<PreluKernel<P>, PreluParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.alpha, para.input]
}
func inferShape() {
// para.output.dim = para.input.dim
}
......
metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift
浏览文件 @ a789f668
......@@ -55,6 +55,11 @@ class PriorBoxParam<P: PrecisionType>: OpParam {
class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input, para.inputImage]
}
func inferShape() {
}
......@@ -66,6 +71,32 @@ class PriorBoxOp<P: PrecisionType>: Operator<PriorBoxKernel<P>, PriorBoxParam<P>
throw error
}
}
func delogOutput() {
print("pribox: ")
print("output: ")
// output
let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in
return o
}
print(outputArray)
// writeToLibrary(fileName: "box_out", array: outputArray)
// output variance
// let outputVarianceArray = para.outputVariances.metalTexture.floatArray { (o: Float32) -> Float32 in
// return o
// }
//
// print(" output variance: \(outputVarianceArray)")
// writeToLibrary(fileName: "variance_out", array: outputVarianceArray)
print("pribox write done ")
}
}
......
metal/paddle-mobile/paddle-mobile/Operators/ReluOp.swift
浏览文件 @ a789f668
......@@ -30,6 +30,10 @@ class ReluParam<P: PrecisionType>: OpParam {
class ReluOp<P: PrecisionType>: Operator<ReluKernel<P>, ReluParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() {
para.output.dim = para.input.dim
}
......
metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift
浏览文件 @ a789f668
......@@ -20,20 +20,24 @@ class ReshapeParam<P: PrecisionType>: OpParam {
do {
input = try ReshapeParam.inputX(inputs: opDesc.inputs, from: inScope)
output = try ReshapeParam.outputOut(outputs: opDesc.outputs, from: inScope)
// shape = output.dim
// shape = output.dim
inplace = try ReshapeParam.getAttr(key: "inplace", attrs: opDesc.attrs)
} catch let error {
throw error
}
}
let input: Texture<P>
// let shape: [Int]
// let shape: [Int]
let inplace: Bool
var output: Texture<P>
}
class ReshapeOp<P: PrecisionType>: Operator<ReshapeKernel<P>, ReshapeParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() {
// para.output.dim = para.input.dim
}
......@@ -48,7 +52,7 @@ class ReshapeOp<P: PrecisionType>: Operator<ReshapeKernel<P>, ReshapeParam<P>>,
}
func delogOutput() {
print("reshape delog")
let _: P? = para.input.metalTexture.logDesc(header: "reshape input: ", stridable: false)
// let _: P? = para.input.metalTexture.logDesc(header: "reshape input: ", stridable: false)
let _: P? = para.output.metalTexture.logDesc(header: "reshape output: ", stridable: false)
}
}
metal/paddle-mobile/paddle-mobile/Operators/SoftmaxOp.swift
浏览文件 @ a789f668
......@@ -30,6 +30,10 @@ class SoftmaxParam<P: PrecisionType>: OpParam {
class SoftmaxOp<P: PrecisionType>: Operator<SoftmaxKernel<P>, SoftmaxParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() {
// para.output.dim = para.input.dim
}
......
metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift
浏览文件 @ a789f668
......@@ -32,8 +32,12 @@ class TransposeParam<P: PrecisionType>: OpParam {
class TransposeOp<P: PrecisionType>: Operator<TransposeKernel<P>, TransposeParam<P>>, Runable, Creator, InferShaperable{
func inputs() -> [Variant] {
return [para.input]
}
func inferShape() {
para.output.dim = para.input.dim
//para.output.dim = para.input.dim
}
typealias OpType = TransposeOp<P>
......@@ -44,6 +48,21 @@ class TransposeOp<P: PrecisionType>: Operator<TransposeKernel<P>, TransposeParam
throw error
}
}
func delogOutput() {
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)
}
}
......
metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift
浏览文件 @ a789f668
......@@ -17,7 +17,7 @@ import Foundation
struct TensorDesc {
let dims: [Int]
let dataType: VarTypeType
let dataLayout: DataLayout = DataLayout.NHWC()
let dataLayout: DataLayout = DataLayout.NCHW()
var NCHWDim: [Int] {
get {
if dims.count != 4 {
......@@ -53,7 +53,7 @@ struct TensorDesc {
}
init(protoTensorDesc: PaddleMobile_Framework_Proto_VarType.TensorDesc) {
dims = protoTensorDesc.dims.map{ Int($0) > 0 ? Int($0) : 1 }
dims = protoTensorDesc.dims.map{ Int($0) > 0 ? Int($0) : abs(Int($0)) }
dataType = VarTypeType.init(rawValue: protoTensorDesc.dataType.rawValue) ?? .ErrorType
}
......
metal/paddle-mobile/paddle-mobile/framework/Tensor.swift
浏览文件 @ a789f668
......@@ -174,7 +174,7 @@ class Tensor<P: PrecisionType>: Tensorial {
fatalError(" not support !")
}
//TODO: release
data.release()
// data.release()
}
var width: Int {
......
metal/paddle-mobile/paddle-mobile/framework/Texture.swift
浏览文件 @ a789f668
......@@ -40,13 +40,14 @@ extension InputTexture {
public class Texture<P: PrecisionType>: Tensorial {
var dim: Dim
var tensorDim: Dim
private(set) public var tensorDim: Dim
private(set) public var originDim: Dim
private var textureDesc: MTLTextureDescriptor!
public var metalTexture: MTLTexture!
var transpose: [Int] = [0, 1, 2, 3]
func initTexture(device: MTLDevice, transpose: [Int] = [0, 1, 2, 3]) {
func initTexture(device: MTLDevice, inTranspose: [Int] = [0, 1, 2, 3]) {
transpose = inTranspose
let newDim = transpose.map { originDim[$0] }
let newLayout = transpose.map { layout.layoutWithDim[$0] }
......@@ -56,10 +57,12 @@ public class Texture<P: PrecisionType>: Tensorial {
let tmpTextureDes = MTLTextureDescriptor.init()
tmpTextureDes.width = layout.W ?? 1
tmpTextureDes.height = layout.H ?? 1
tmpTextureDes.width = newDim[2]
// layout.W ?? 1
tmpTextureDes.height = newDim[1]
// layout.H ?? 1
tmpTextureDes.depth = 1
tmpTextureDes.arrayLength = ((layout.N ?? 1) * (layout.C ?? 1) + 3) / 4
tmpTextureDes.arrayLength = ((newDim[0]) * (newDim[3]) + 3) / 4
tmpTextureDes.textureType = .type2DArray
if MemoryLayout<P>.size == 1 {
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或