未验证 提交 b602b048 编写于 作者: R Ruilong Liu 提交者: GitHub

Merge pull request #837 from codeWorm2015/metal

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