Merge pull request #1444 from codeWorm2015/develop

use metal lib project manager metal files

metal/MobileNetDemo/MobileNetDemo.xcodeproj/project.pbxproj

@@ -15,6 +15,7 @@
 		FCB40E5121E0CEBB0075EC91 /* mobilenet_model in Resources */ = {isa = PBXBuildFile; fileRef = FCB40E4F21E0CEBB0075EC91 /* mobilenet_model */; };
 		FCB40E5221E0CEBB0075EC91 /* mobilenet_params in Resources */ = {isa = PBXBuildFile; fileRef = FCB40E5021E0CEBB0075EC91 /* mobilenet_params */; };
 		FCB40E5421E0CEF80075EC91 /* synset.txt in Resources */ = {isa = PBXBuildFile; fileRef = FCB40E5321E0CEF80075EC91 /* synset.txt */; };
+		FCC15E13221E715400DC3CB2 /* paddle-mobile-metallib.metallib in Resources */ = {isa = PBXBuildFile; fileRef = FCC15E12221E715400DC3CB2 /* paddle-mobile-metallib.metallib */; };
 		FCD3873821E1C31F0052F3D0 /* paddle_mobile.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = FCD3873721E1C31F0052F3D0 /* paddle_mobile.framework */; };
 		FCD3873921E1C31F0052F3D0 /* paddle_mobile.framework in Embed Frameworks */ = {isa = PBXBuildFile; fileRef = FCD3873721E1C31F0052F3D0 /* paddle_mobile.framework */; settings = {ATTRIBUTES = (CodeSignOnCopy, RemoveHeadersOnCopy, ); }; };
 		FCF2870921DFAEC7009A87DA /* AppDelegate.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCF2870821DFAEC7009A87DA /* AppDelegate.swift */; };
@@ -49,6 +50,7 @@
 		FCB40E4F21E0CEBB0075EC91 /* mobilenet_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = mobilenet_model; sourceTree = "<group>"; };
 		FCB40E5021E0CEBB0075EC91 /* mobilenet_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = mobilenet_params; sourceTree = "<group>"; };
 		FCB40E5321E0CEF80075EC91 /* synset.txt */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = synset.txt; sourceTree = "<group>"; };
+		FCC15E12221E715400DC3CB2 /* paddle-mobile-metallib.metallib */ = {isa = PBXFileReference; lastKnownFileType = "archive.metal-library"; name = "paddle-mobile-metallib.metallib"; path = "../../../../Library/Developer/Xcode/DerivedData/paddle-mobile-hdsimtkoxoondndnjczkbkchcwyh/Build/Products/Release-iphoneos/paddle-mobile-metallib.metallib"; sourceTree = "<group>"; };
 		FCD3873721E1C31F0052F3D0 /* paddle_mobile.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; path = paddle_mobile.framework; sourceTree = BUILT_PRODUCTS_DIR; };
 		FCF2870521DFAEC7009A87DA /* MobileNetDemo.app */ = {isa = PBXFileReference; explicitFileType = wrapper.application; includeInIndex = 0; path = MobileNetDemo.app; sourceTree = BUILT_PRODUCTS_DIR; };
 		FCF2870821DFAEC7009A87DA /* AppDelegate.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = AppDelegate.swift; sourceTree = "<group>"; };
@@ -127,6 +129,7 @@
 		FCF286FC21DFAEC7009A87DA = {
 			isa = PBXGroup;
 			children = (
+				FCC15E12221E715400DC3CB2 /* paddle-mobile-metallib.metallib */,
 				FCD3873721E1C31F0052F3D0 /* paddle_mobile.framework */,
 				FCF2870721DFAEC7009A87DA /* MobileNetDemo */,
 				FCF2870621DFAEC7009A87DA /* Products */,
@@ -225,6 +228,7 @@
 				FCB40E5121E0CEBB0075EC91 /* mobilenet_model in Resources */,
 				FCB40DE921E0B9410075EC91 /* banana.jpeg in Resources */,
 				FCF2871021DFAEC8009A87DA /* Assets.xcassets in Resources */,
+				FCC15E13221E715400DC3CB2 /* paddle-mobile-metallib.metallib in Resources */,
 				FCB40E5421E0CEF80075EC91 /* synset.txt in Resources */,
 				FCB40E5221E0CEBB0075EC91 /* mobilenet_params in Resources */,
 				FCF2870E21DFAEC7009A87DA /* Main.storyboard in Resources */,

metal/MobileNetDemo/MobileNetDemo/MobileNet.swift

@@ -46,7 +46,7 @@ public class MobileNet: Net{
     let firstRes = res[0]
     let resPointer = firstRes.result
     var s: [String] = []
-    (0..<res.capacity).map { resPointer[$0] }.top(r: 5).enumerated().forEach{
+    (0..<firstRes.capacity).map { resPointer[$0] }.top(r: 5).enumerated().forEach{
       s.append(String(format: "%d: %@ (%3.2f%%)", $0 + 1, labels[$1.0], $1.1 * 100))
     }
     return s.joined(separator: "\n")
@@ -59,6 +59,8 @@ public class MobileNet: Net{
     paramPath = Bundle.main.path(forResource: "mobilenet_params", ofType: nil) ?! "para null"
     preprocessKernel = MobilenetPreProccess.init(device: device)
     inputDim = Dim.init(inDim: [1, 224, 224, 3])
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
   }
 }
 

metal/Podfile

@@ -27,3 +27,8 @@ target 'MobileNetDemo' do
     pod 'Protobuf', '~> 3.0.0'
 end
 
+
+target 'paddle-mobile-metallib' do
+    project 'paddle-mobile-metallib/paddle-mobile-metallib.xcodeproj'
+end
+

metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj

@@ -45,6 +45,7 @@
 		FC9797C721D609FB00F2FD90 /* synset.txt in Resources */ = {isa = PBXBuildFile; fileRef = FC9797C621D609FB00F2FD90 /* synset.txt */; };
 		FC9797CF21D6506F00F2FD90 /* mingren.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FC9797CE21D6506F00F2FD90 /* mingren.jpg */; };
 		FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC542122EF5400D94F7E /* MetalHelper.swift */; };
+		FCC15E15221E716500DC3CB2 /* paddle-mobile-metallib.metallib in Resources */ = {isa = PBXBuildFile; fileRef = FCC15E14221E716400DC3CB2 /* paddle-mobile-metallib.metallib */; };
 		FCCED60521D7646E00BE8D5F /* test_image_super in Resources */ = {isa = PBXBuildFile; fileRef = FCCED60421D7646E00BE8D5F /* test_image_super */; };
 		FCEBEC2C20E1391F00C0B14D /* paddle_mobile.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */; };
 		FCEBEC2D20E1391F00C0B14D /* paddle_mobile.framework in Embed Frameworks */ = {isa = PBXBuildFile; fileRef = FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */; settings = {ATTRIBUTES = (CodeSignOnCopy, RemoveHeadersOnCopy, ); }; };
@@ -113,6 +114,7 @@
 		FC9797C621D609FB00F2FD90 /* synset.txt */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = synset.txt; sourceTree = "<group>"; };
 		FC9797CE21D6506F00F2FD90 /* mingren.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = mingren.jpg; sourceTree = "<group>"; };
 		FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = "<group>"; };
+		FCC15E14221E716400DC3CB2 /* paddle-mobile-metallib.metallib */ = {isa = PBXFileReference; lastKnownFileType = "archive.metal-library"; name = "paddle-mobile-metallib.metallib"; path = "../../../../Library/Developer/Xcode/DerivedData/paddle-mobile-hdsimtkoxoondndnjczkbkchcwyh/Build/Products/Release-iphoneos/paddle-mobile-metallib.metallib"; sourceTree = "<group>"; };
 		FCCED60421D7646E00BE8D5F /* test_image_super */ = {isa = PBXFileReference; lastKnownFileType = file; path = test_image_super; sourceTree = "<group>"; };
 		FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; path = paddle_mobile.framework; sourceTree = BUILT_PRODUCTS_DIR; };
 		FCF437E7214B6DDB00943429 /* MultiPredictViewController.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MultiPredictViewController.swift; sourceTree = "<group>"; };
@@ -153,6 +155,7 @@
 		FC039B7520E11C550081E9F8 = {
 			isa = PBXGroup;
 			children = (
+				FCC15E14221E716400DC3CB2 /* paddle-mobile-metallib.metallib */,
 				FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */,
 				FC039B8020E11C550081E9F8 /* paddle-mobile-demo */,
 				FC039B7F20E11C550081E9F8 /* Products */,
@@ -392,6 +395,7 @@
 				FC9797C221D608E000F2FD90 /* mobilenet_model in Resources */,
 				FC2BFCC021DF279900C262B2 /* classify-img-output.png in Resources */,
 				FC203FB221CBFDBA00B37166 /* test.jpg in Resources */,
+				FCC15E15221E716500DC3CB2 /* paddle-mobile-metallib.metallib in Resources */,
 				FC704C2321D237FC00F98BAB /* combined_mobilenet_model in Resources */,
 				FC9797C321D608E000F2FD90 /* mobilenet_params in Resources */,
 				FC704C2421D237FC00F98BAB /* yolo_params in Resources */,

metal/paddle-mobile-demo/paddle-mobile-demo/Base.lproj/Main.storyboard

 <?xml version="1.0" encoding="UTF-8"?>
-<document type="com.apple.InterfaceBuilder3.CocoaTouch.Storyboard.XIB" version="3.0" toolsVersion="14460.31" targetRuntime="iOS.CocoaTouch" propertyAccessControl="none" useAutolayout="YES" useTraitCollections="YES" useSafeAreas="YES" colorMatched="YES" initialViewController="BYZ-38-t0r">
+<document type="com.apple.InterfaceBuilder3.CocoaTouch.Storyboard.XIB" version="3.0" toolsVersion="14460.31" targetRuntime="iOS.CocoaTouch" propertyAccessControl="none" useAutolayout="YES" useTraitCollections="YES" useSafeAreas="YES" colorMatched="YES" initialViewController="4MS-jc-i6A">
     <device id="retina4_7" orientation="portrait">
         <adaptation id="fullscreen"/>
     </device>
@@ -246,10 +246,6 @@
                         <rect key="frame" x="0.0" y="0.0" width="375" height="667"/>
                         <autoresizingMask key="autoresizingMask" widthSizable="YES" heightSizable="YES"/>
                         <subviews>
-                            <imageView userInteractionEnabled="NO" contentMode="scaleToFill" horizontalHuggingPriority="251" verticalHuggingPriority="251" fixedFrame="YES" translatesAutoresizingMaskIntoConstraints="NO" id="2p5-S3-M4T">
-                                <rect key="frame" x="16" y="63" width="240" height="128"/>
-                                <autoresizingMask key="autoresizingMask" flexibleMaxX="YES" flexibleMaxY="YES"/>
-                            </imageView>
                             <button opaque="NO" contentMode="scaleToFill" contentHorizontalAlignment="center" contentVerticalAlignment="center" buttonType="roundedRect" lineBreakMode="middleTruncation" translatesAutoresizingMaskIntoConstraints="NO" id="37q-nm-0H7">
                                 <rect key="frame" x="38" y="610" width="42" height="30"/>
                                 <constraints>
@@ -287,7 +283,16 @@
                                     <constraint firstAttribute="height" constant="30" id="eAt-Uc-BxX"/>
                                 </constraints>
                                 <state key="normal" title="clear"/>
+                                <connections>
+                                    <action selector="clear:" destination="4MS-jc-i6A" eventType="touchUpInside" id="yW8-Dq-qwU"/>
+                                </connections>
                             </button>
+                            <imageView userInteractionEnabled="NO" contentMode="scaleToFill" horizontalHuggingPriority="251" verticalHuggingPriority="251" translatesAutoresizingMaskIntoConstraints="NO" id="2p5-S3-M4T">
+                                <rect key="frame" x="0.0" y="20" width="375" height="211"/>
+                                <constraints>
+                                    <constraint firstAttribute="width" secondItem="2p5-S3-M4T" secondAttribute="height" multiplier="16:9" id="9Gh-8L-t3g"/>
+                                </constraints>
+                            </imageView>
                         </subviews>
                         <color key="backgroundColor" white="1" alpha="1" colorSpace="custom" customColorSpace="genericGamma22GrayColorSpace"/>
                         <constraints>
@@ -295,17 +300,23 @@
                             <constraint firstItem="DZa-sd-lY7" firstAttribute="leading" secondItem="pdS-6e-Pd1" secondAttribute="trailing" constant="45" id="8dB-uI-cs9"/>
                             <constraint firstItem="fAg-ai-yaA" firstAttribute="leading" secondItem="37q-nm-0H7" secondAttribute="trailing" constant="39" id="EAV-Oq-jeD"/>
                             <constraint firstItem="vsb-FH-h7h" firstAttribute="bottom" secondItem="fAg-ai-yaA" secondAttribute="bottom" constant="27" id="Px0-A9-Eql"/>
+                            <constraint firstItem="2p5-S3-M4T" firstAttribute="leading" secondItem="vsb-FH-h7h" secondAttribute="leading" id="RNx-6D-oix"/>
                             <constraint firstItem="pdS-6e-Pd1" firstAttribute="leading" secondItem="fAg-ai-yaA" secondAttribute="trailing" constant="32" id="ZUR-Nv-aNb"/>
+                            <constraint firstItem="2p5-S3-M4T" firstAttribute="top" secondItem="vsb-FH-h7h" secondAttribute="top" id="atk-ma-aSA"/>
                             <constraint firstItem="vsb-FH-h7h" firstAttribute="bottom" secondItem="pdS-6e-Pd1" secondAttribute="bottom" constant="27" id="kPx-mt-ab9"/>
+                            <constraint firstItem="2p5-S3-M4T" firstAttribute="trailing" secondItem="vsb-FH-h7h" secondAttribute="trailing" id="mwX-bu-jJY"/>
                             <constraint firstItem="37q-nm-0H7" firstAttribute="leading" secondItem="vsb-FH-h7h" secondAttribute="leading" constant="38" id="trH-Fq-sSv"/>
                             <constraint firstItem="vsb-FH-h7h" firstAttribute="bottom" secondItem="DZa-sd-lY7" secondAttribute="bottom" constant="27" id="yNJ-hq-2Qg"/>
                         </constraints>
                         <viewLayoutGuide key="safeArea" id="vsb-FH-h7h"/>
                     </view>
+                    <connections>
+                        <outlet property="imageView" destination="2p5-S3-M4T" id="ePO-1L-eb4"/>
+                    </connections>
                 </viewController>
                 <placeholder placeholderIdentifier="IBFirstResponder" id="hGb-Pb-icS" userLabel="First Responder" sceneMemberID="firstResponder"/>
             </objects>
-            <point key="canvasLocation" x="-721" y="-427"/>
+            <point key="canvasLocation" x="-135.19999999999999" y="-218.1409295352324"/>
         </scene>
     </scenes>
     <resources>

metal/paddle-mobile-demo/paddle-mobile-demo/MetalHelper.swift

@@ -17,14 +17,15 @@ import MetalKit
 import Foundation
 import paddle_mobile
 
-public class MetalHelper {
-  let device: MTLDevice
-  let queue: MTLCommandQueue
-  let textureLoader: MTKTextureLoader
-  static let shared: MetalHelper = MetalHelper.init()
-  private init(){
+@objc public class MetalHelper: NSObject {
+  @objc let device: MTLDevice
+  @objc let queue: MTLCommandQueue
+  @objc let textureLoader: MTKTextureLoader
+  @objc static let shared: MetalHelper = MetalHelper.init()
+  private override init(){
     device = MTLCreateSystemDefaultDevice()!
     queue = device.makeCommandQueue()!
     textureLoader = MTKTextureLoader.init(device: device)
+    super.init()
   }
 }

metal/paddle-mobile-demo/paddle-mobile-demo/Net/Genet.swift

@@ -22,14 +22,18 @@ public class Genet: Net {
     paramPath = Bundle.main.path(forResource: "genet_params", ofType: nil) ?! "para null"
     preprocessKernel = GenetPreProccess.init(device: device)
     inputDim = Dim.init(inDim: [1, 128, 128, 3])
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
   }
   
-  @objc override public init(device: MTLDevice, paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
+  @objc override public init(device: MTLDevice, inParamPointer: UnsafeMutableRawPointer, inParamSize:Int, inModelPointer: UnsafeMutableRawPointer, inModelSize: Int) {
     super.init(device: device,
-               paramPointer: paramPointer,
-               paramSize: paramSize,
-               modePointer: modePointer,
-               modelSize: modelSize)
+               inParamPointer: inParamPointer,
+               inParamSize: inParamSize,
+               inModelPointer: inModelPointer,
+               inModelSize: inModelSize)
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
     preprocessKernel = GenetPreProccess.init(device: device)
     inputDim = Dim.init(inDim: [1, 128, 128, 3])
   }

metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNet.swift

@@ -46,7 +46,7 @@ public class MobileNet: Net{
   override public func resultStr(res: [ResultHolder]) -> String {
     let resPointer = res[0].result
     var s: [String] = []
-    (0..<res.capacity).map { resPointer[$0] }.top(r: 5).enumerated().forEach{
+    (0..<res[0].capacity).map { resPointer[$0] }.top(r: 5).enumerated().forEach{
       s.append(String(format: "%d: %@ (%3.2f%%)", $0 + 1, labels[$1.0], $1.1 * 100))
     }
     return s.joined(separator: "\n")
@@ -61,6 +61,8 @@ public class MobileNet: Net{
 //    metalLibPath = Bundle.main.path(forResource: "PaddleMobileMetal", ofType: "metallib") ?! " can't be nil "
     preprocessKernel = MobilenetPreProccess.init(device: device)
     inputDim = Dim.init(inDim: [1, 224, 224, 3])
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
   }
 }
 

metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNetCombined.swift

@@ -22,8 +22,8 @@ public class MobileNetCombined: Net {
     modelPath = Bundle.main.path(forResource: "combined_mobilenet_model", ofType: nil) ?! "model null"
     paramPath = Bundle.main.path(forResource: "combined_mobilenet_params", ofType: nil) ?! "para null"
     inputDim = Dim.init(inDim: [1, 224, 224, 3])
-//    metalLoadMode = .LoadMetalInCustomMetalLib
-//    metalLibPath = Bundle.main.path(forResource: "PaddleMobileMetal", ofType: "metallib") ?! " can't be nil "
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
   }
   
   override  public func resultStr(res: [ResultHolder]) -> String {

metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNetSSD.swift

@@ -21,19 +21,19 @@ public class MobileNet_ssd_hand: Net {
     except = 2
     modelPath = Bundle.main.path(forResource: "ssd_hand_model", ofType: nil) ?! "model null"
     paramPath = Bundle.main.path(forResource: "ssd_hand_params", ofType: nil) ?! "para null"
-//    metalLoadMode = .LoadMetalInCustomMetalLib
-//    metalLibPath = Bundle.main.path(forResource: "PaddleMobileMetal", ofType: "metallib") ?! " can't be nil "
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
     preprocessKernel = MobilenetssdPreProccess.init(device: device)
     inputDim = Dim.init(inDim: [1, 300, 300, 3])
   }
   
-  @objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
-    super.init(device:device,paramPointer:paramPointer,paramSize:paramSize,modePointer:modePointer,modelSize:modelSize)
+  @objc override public init(device: MTLDevice,inParamPointer: UnsafeMutableRawPointer, inParamSize:Int, inModelPointer inModePointer: UnsafeMutableRawPointer, inModelSize: Int) {
+    super.init(device:device,inParamPointer:inParamPointer,inParamSize:inParamSize,inModelPointer:inModePointer,inModelSize:inModelSize)
     except = 2
     modelPath = ""
     paramPath = ""
-//    metalLoadMode = .LoadMetalInCustomMetalLib
-//    metalLibPath = Bundle.main.path(forResource: "PaddleMobileMetal", ofType: "metallib") ?! " can't be nil "
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
     preprocessKernel = MobilenetssdPreProccess.init(device: device)
     inputDim = Dim.init(inDim: [1, 300, 300, 3])
   }

metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobilenetSSD_AR.swift

@@ -23,13 +23,17 @@ public class MobileNet_ssd_AR: Net {
     paramPath = Bundle.main.path(forResource: "ar_params", ofType: nil) ?! "para null"
     preprocessKernel = MobilenetssdPreProccess.init(device: device)
     inputDim = Dim.init(inDim: [1, 160, 160, 3])
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
   }
   
-  @objc override public init(device: MTLDevice,paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
-    super.init(device:device,paramPointer:paramPointer,paramSize:paramSize,modePointer:modePointer,modelSize:modelSize)
+  @objc override public init(device: MTLDevice, inParamPointer: UnsafeMutableRawPointer, inParamSize:Int, inModelPointer: UnsafeMutableRawPointer, inModelSize: Int) {
+    super.init(device:device,inParamPointer:inParamPointer,inParamSize:inParamSize,inModelPointer:inModelPointer,inModelSize:inModelSize)
     except = 2
     preprocessKernel = MobilenetssdPreProccess.init(device: device)
     inputDim = Dim.init(inDim: [1, 160, 160, 3])
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
   }
   
   class MobilenetssdPreProccess: CusomKernel {

metal/paddle-mobile-demo/paddle-mobile-demo/Net/YoloNet.swift

@@ -23,8 +23,8 @@ public class YoloNet: Net {
     modelPath = Bundle.main.path(forResource: "yolo_model", ofType: nil) ?! "model null"
     paramPath = Bundle.main.path(forResource: "yolo_params", ofType: nil) ?! "para null"
     inputDim = Dim.init(inDim: [1, 416, 416, 3])
-//    metalLoadMode = .LoadMetalInCustomMetalLib
-//    metalLibPath = Bundle.main.path(forResource: "PaddleMobileMetal", ofType: "metallib") ?! " can't be nil "
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
   }
 
   override  public func resultStr(res: [ResultHolder]) -> String {

metal/paddle-mobile-demo/paddle-mobile-demo/OCDemo/LoadPointerViewController.h

@@ -14,6 +14,10 @@
 
 #import <UIKit/UIKit.h>
 
+
+/**
+ @b 从内存中加载模型 Demo, 可以在 main storyboard 中调整 Demo
+ */
 @interface LoadPointerViewController : UIViewController
 
 @end

metal/paddle-mobile-demo/paddle-mobile-demo/OCDemo/LoadPointerViewController.m

@@ -13,17 +13,20 @@
  limitations under the License. */
 
 #import "PaddleMobileGPU.h"
+#import "paddle_mobile_demo-Swift.h"
 #import "LoadPointerViewController.h"
-#import "paddle-mobile-demo-Bridging-Header.h"
 
 #import <Metal/Metal.h>
+#import <MetalKit/MetalKit.h>
 
 @interface LoadPointerViewController ()
 
-@property (strong, nonatomic) id<MTLDevice> device;
+@property (weak, nonatomic) IBOutlet UIImageView *imageView;
+
+@property (assign, nonatomic) BOOL loaded;
 @property (strong, nonatomic) id<MTLTexture> texture;
-@property (strong, nonatomic) id<MTLCommandQueue> queue;
-@property (strong, nonatomic) PaddleMobileGPU *runner;
+
+@property (strong, nonatomic) PaddleMobileGPU *paddleMobile;
 @property (strong, nonatomic) ModelConfig *modelConfig;
 
 @end
@@ -31,19 +34,12 @@
 @implementation LoadPointerViewController
 
 - (void)viewDidLoad {
-    [super viewDidLoad];
-  
-  
-  self.device = MTLCreateSystemDefaultDevice();
+  [super viewDidLoad];
   
-  self.queue = [self.device newCommandQueue];
+  self.imageView.image = [UIImage imageNamed:@"banana.jpeg"];
   
-    // Do any additional setup after loading the view.
-//  NSString *modelPath = [[NSBundle mainBundle] URLForResource:@"genet_model" withExtension:nil].path;
-//  NSString *paramPath = [[NSBundle mainBundle] URLForResource:@"genet_params" withExtension:nil].path;
-  
-  NSString *modelPath = [[NSBundle mainBundle] URLForResource:@"ar_model" withExtension:nil].path;
-  NSString *paramPath = [[NSBundle mainBundle] URLForResource:@"ar_params" withExtension:nil].path;
+  NSString *modelPath = [[NSBundle mainBundle] URLForResource:@"super_model" withExtension:nil].path;
+  NSString *paramPath = [[NSBundle mainBundle] URLForResource:@"super_params" withExtension:nil].path;
 
   long fileSize;
   FILE *fp;
@@ -66,114 +62,55 @@
   fclose(parmaFilePointer);
   
   _modelConfig = [[ModelConfig alloc] init];
-//  _modelConfig.means = @[[NSNumber numberWithFloat:128.0], [NSNumber numberWithFloat:128.0], [NSNumber numberWithFloat:128.0]];
-//  _modelConfig.scale = 0.017;
-//  _modelConfig.dims = @[[NSNumber numberWithFloat:1], [NSNumber numberWithFloat:128.], [NSNumber numberWithFloat:128.0],[NSNumber numberWithFloat:3.0]];
-  _modelConfig.means = @[[NSNumber numberWithFloat:103.94], [NSNumber numberWithFloat:116.78], [NSNumber numberWithFloat:123.68]];
-  _modelConfig.scale = 1;
-  _modelConfig.dims = @[[NSNumber numberWithFloat:1], [NSNumber numberWithFloat:160.], [NSNumber numberWithFloat:160.0],[NSNumber numberWithFloat:3.0]];
   _modelConfig.modelPointer = buffer;
   _modelConfig.modelSize = (int)fileSize;
   _modelConfig.paramPointer = parmaBuffer;
   _modelConfig.paramSize = (int)paramfileSize;
 }
 - (IBAction)loaderButtonPressed:(id)sender {
-//  _runner = [[PaddleMobileGPU alloc] initWithCommandQueue:self.queue net:GenetType modelConfig:_modelConfig];
-  _runner = [[PaddleMobileGPU alloc] initWithCommandQueue:self.queue net:MobileNetSSDType modelConfig:_modelConfig];
-  
-  [_runner load];
+  self.paddleMobile = [[PaddleMobileGPU alloc] initWithCommandQueue:MetalHelper.shared.queue net:SuperResolutionNetType modelConfig:_modelConfig];
+  _loaded = [self.paddleMobile load];
+  NSLog(@" load 结果: %@", _loaded ? @"成功" : @"失败");
 }
 - (IBAction)predictButtonPressed:(id)sender {
   [self predict];
 }
 
-- (id<MTLTexture>) createTextureFromImage:(UIImage*) image device:(id<MTLDevice>) device
-{
-  image  =[UIImage imageWithCGImage:[image CGImage]
-                              scale:[image scale]
-                        orientation: UIImageOrientationLeft];
-  
-  NSLog(@"orientation and size and stuff %ld %f %f", (long)image.imageOrientation, image.size.width, image.size.height);
-  
-  CGImageRef imageRef = image.CGImage;
-  
-  size_t width = self.view.frame.size.width;
-  size_t height = self.view.frame.size.height;
-  
-  size_t bitsPerComponent = CGImageGetBitsPerComponent(imageRef);
-  size_t bitsPerPixel = CGImageGetBitsPerPixel(imageRef);
-  
-  CGColorSpaceRef colorSpace = CGImageGetColorSpace(imageRef);
-  
-  CGImageAlphaInfo alphaInfo = CGImageGetAlphaInfo(imageRef);
-  
-  //  NSLog(@"%@ %u", colorSpace, alphaInfo);
-  
-  CGBitmapInfo bitmapInfo = kCGBitmapByteOrderDefault | alphaInfo;
-  //    NSLog(@"bitmap info %u", bitmapInfo);
-  
-  
-  CGContextRef context = CGBitmapContextCreate( NULL, width, height, bitsPerComponent, (bitsPerPixel / 8) * width, colorSpace, bitmapInfo);
-  
-  if( !context )
-  {
-    NSLog(@"Failed to load image, probably an unsupported texture type");
-    return nil;
+- (void)predict {
+  UIImage *image = self.imageView.image;
+  if (!image) {
+    NSLog(@" image is nil");
+    return;
+  }
+  id<MTLTexture> texture = [MetalHelper.shared.textureLoader newTextureWithCGImage:image.CGImage options:nil error:nil];
+  _texture = texture;
+  if (!_texture) {
+    NSLog(@" texture is nil");
+    return;
   }
   
-  CGContextDrawImage( context, CGRectMake( 0, 0, width, height ), image.CGImage);
-  
-  
-  MTLPixelFormat format = MTLPixelFormatRGBA8Unorm;
-  
-  MTLTextureDescriptor *texDesc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:format
-                                                                                     width:width
-                                                                                    height:height
-                                                                                 mipmapped:NO];
-  id<MTLTexture> texture = [device newTextureWithDescriptor:texDesc];
-  
-  [texture replaceRegion:MTLRegionMake2D(0, 0, width, height)
-             mipmapLevel:0
-               withBytes:CGBitmapContextGetData(context)
-             bytesPerRow:4 * width];
+  if (!self.loaded) {
+    NSLog(@" not load ");
+    return;
+  }
   
-  return texture;
-}
-
-- (void)predict {
-  _texture = [self createTextureFromImage:[UIImage imageNamed:@"hand.jpg"] device:self.device];
   NSTimeInterval startTime = [[NSDate date] timeIntervalSince1970];
-  NSInteger max = 428;
+  NSInteger max = 1;
   for (int i = 0;i < max; i ++) {
-    [_runner predict:_texture withCompletion:^(BOOL success , NSArray<NSNumber *> *result) {
+    [self.paddleMobile predict:_texture withCompletion:^(BOOL success , NSArray<NSNumber *> *result) {
       if (success) {
         if (i == max -1) {
           double time = [[NSDate date] timeIntervalSince1970] - startTime;
           time = (time/max)*1000;
           NSLog(@"gap ==== %fms",time);
         }
-//        for (int i = 0; i < result.count; i ++) {
-//          NSNumber *number = result[i];
-//          NSLog(@"result %d = %f:",i, [number floatValue]);
-//        }
       }
     }];
   }
 }
-
-- (void)didReceiveMemoryWarning {
-    [super didReceiveMemoryWarning];
-    // Dispose of any resources that can be recreated.
-}
-
-/*
-#pragma mark - Navigation
-
-// In a storyboard-based application, you will often want to do a little preparation before navigation
-- (void)prepareForSegue:(UIStoryboardSegue *)segue sender:(id)sender {
-    // Get the new view controller using [segue destinationViewController].
-    // Pass the selected object to the new view controller.
+- (IBAction)clear:(id)sender {
+  [self.paddleMobile clear];
+  self.loaded = NO;
 }
-*/
 
 @end

metal/paddle-mobile-demo/paddle-mobile-demo/OCInterface/PaddleMobileGPU.m

@@ -53,9 +53,9 @@
   if (self) {
     Net *net = nil;
     if (netType == SuperResolutionNetType) {
-      net = [[SuperResolutionNet alloc] initWithDevice:queue.device];
+      net = [[SuperResolutionNet alloc] initWithDevice:queue.device inParamPointer:config.paramPointer inParamSize:config.paramSize inModelPointer:config.modelPointer inModelSize:config.modelSize];
     } else if (netType == MobileNetSSDType) {
-      net = [[MobileNet_ssd_AR alloc] initWithDevice:queue.device paramPointer:config.paramPointer paramSize:config.paramSize modePointer:config.modelPointer modelSize:config.modelSize];
+      net = [[MobileNet_ssd_AR alloc] initWithDevice:queue.device inParamPointer:config.paramPointer inParamSize:config.paramSize inModelPointer:config.modelPointer inModelSize:config.modelSize];
     }
     runner = [[Runner alloc] initInNet:net commandQueue:queue];
   }

metal/paddle-mobile-demo/paddle-mobile-demo/OCInterface/SuperResolutionNet.swift

@@ -20,6 +20,18 @@ import paddle_mobile
     return "未实现"
   }
   
+  public override init(device: MTLDevice, inParamPointer: UnsafeMutableRawPointer, inParamSize: Int, inModelPointer: UnsafeMutableRawPointer, inModelSize: Int) {
+    super.init(device: device)
+    except = 0
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
+    inputDim = Dim.init(inDim: [1, 224, 224, 3])
+    self.paramPointer = inParamPointer
+    self.paramSize = inParamSize
+    self.modelPointer = inModelPointer
+    self.modelSize = inModelSize
+  }
+    
   @objc override public init(device: MTLDevice) {
     super.init(device: device)
     except = 0
@@ -27,8 +39,8 @@ import paddle_mobile
     paramPath = Bundle.main.path(forResource: "super_params", ofType: nil) ?! "para null"
     preprocessKernel = nil
     inputDim = Dim.init(inDim: [1, 224, 224, 1])
-//    metalLoadMode = .LoadMetalInCustomMetalLib
-//    metalLibPath = Bundle.main.path(forResource: "PaddleMobileMetal", ofType: "metallib") ?! " can't be nil "
+    metalLoadMode = .LoadMetalInCustomMetalLib
+    metalLibPath = Bundle.main.path(forResource: "paddle-mobile-metallib", ofType: "metallib")
   }
   
   override public func updateProgram(program: Program) {

metal/paddle-mobile-metallib/paddle-mobile-metallib.xcodeproj/project.pbxproj

+// !$*UTF8*$!
+{
+	archiveVersion = 1;
+	classes = {
+	};
+	objectVersion = 50;
+	objects = {
+
+/* Begin PBXBuildFile section */
+		5CCC0CF6759710BAFE999DB7 /* Pods_paddle_mobile_metallib.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 5D9D330A035906298947080B /* Pods_paddle_mobile_metallib.framework */; };
+		FCC15DE5221E69E100DC3CB2 /* ReluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DBC221E69DD00DC3CB2 /* ReluKernel.metal */; };
+		FCC15DE6221E69E100DC3CB2 /* BoxCoder.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DBD221E69DD00DC3CB2 /* BoxCoder.metal */; };
+		FCC15DE7221E69E100DC3CB2 /* ConvAddBNReluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DBE221E69DD00DC3CB2 /* ConvAddBNReluKernel.metal */; };
+		FCC15DE8221E69E100DC3CB2 /* Split.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DBF221E69DD00DC3CB2 /* Split.metal */; };
+		FCC15DE9221E69E100DC3CB2 /* BilinearInterp.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC0221E69DD00DC3CB2 /* BilinearInterp.metal */; };
+		FCC15DEA221E69E100DC3CB2 /* ElementwiseAddPreluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC1221E69DD00DC3CB2 /* ElementwiseAddPreluKernel.metal */; };
+		FCC15DEB221E69E100DC3CB2 /* NMSFetchResultKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC2221E69DD00DC3CB2 /* NMSFetchResultKernel.metal */; };
+		FCC15DEC221E69E100DC3CB2 /* Softmax.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC3221E69DD00DC3CB2 /* Softmax.metal */; };
+		FCC15DED221E69E100DC3CB2 /* PoolKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC4221E69DE00DC3CB2 /* PoolKernel.inc.metal */; };
+		FCC15DEE221E69E100DC3CB2 /* ConvTransposeKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC5221E69DE00DC3CB2 /* ConvTransposeKernel.metal */; };
+		FCC15DEF221E69E100DC3CB2 /* Macro.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC6221E69DE00DC3CB2 /* Macro.metal */; };
+		FCC15DF0221E69E100DC3CB2 /* PreluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC7221E69DE00DC3CB2 /* PreluKernel.metal */; };
+		FCC15DF1221E69E100DC3CB2 /* BilinearInterp.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC8221E69DE00DC3CB2 /* BilinearInterp.inc.metal */; };
+		FCC15DF2221E69E100DC3CB2 /* TransposeKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DC9221E69DE00DC3CB2 /* TransposeKernel.inc.metal */; };
+		FCC15DF3221E69E100DC3CB2 /* ConcatKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DCA221E69DE00DC3CB2 /* ConcatKernel.metal */; };
+		FCC15DF4221E69E100DC3CB2 /* ResizeBilinear.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DCB221E69DE00DC3CB2 /* ResizeBilinear.metal */; };
+		FCC15DF5221E69E100DC3CB2 /* Common.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DCC221E69DE00DC3CB2 /* Common.metal */; };
+		FCC15DF6221E69E100DC3CB2 /* PoolKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DCD221E69DE00DC3CB2 /* PoolKernel.metal */; };
+		FCC15DF7221E69E100DC3CB2 /* ReshapeKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DCE221E69DE00DC3CB2 /* ReshapeKernel.inc.metal */; };
+		FCC15DF8221E69E100DC3CB2 /* ConvBNReluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DCF221E69DE00DC3CB2 /* ConvBNReluKernel.metal */; };
+		FCC15DF9221E69E100DC3CB2 /* Kernels.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD0221E69DE00DC3CB2 /* Kernels.metal */; };
+		FCC15DFA221E69E100DC3CB2 /* Shape.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD1221E69DF00DC3CB2 /* Shape.metal */; };
+		FCC15DFB221E69E100DC3CB2 /* Softmax.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD2221E69DF00DC3CB2 /* Softmax.inc.metal */; };
+		FCC15DFC221E69E100DC3CB2 /* ConvAddPreluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD3221E69DF00DC3CB2 /* ConvAddPreluKernel.metal */; };
+		FCC15DFD221E69E100DC3CB2 /* Elementwise.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD4221E69DF00DC3CB2 /* Elementwise.metal */; };
+		FCC15DFE221E69E100DC3CB2 /* ReshapeKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD5221E69DF00DC3CB2 /* ReshapeKernel.metal */; };
+		FCC15DFF221E69E100DC3CB2 /* Scale.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD6221E69DF00DC3CB2 /* Scale.metal */; };
+		FCC15E00221E69E100DC3CB2 /* ConvKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD7221E69DF00DC3CB2 /* ConvKernel.metal */; };
+		FCC15E01221E69E100DC3CB2 /* PriorBoxKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD8221E69DF00DC3CB2 /* PriorBoxKernel.metal */; };
+		FCC15E02221E69E100DC3CB2 /* BatchNormRelu.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DD9221E69E000DC3CB2 /* BatchNormRelu.metal */; };
+		FCC15E03221E69E100DC3CB2 /* TransposeKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DDA221E69E000DC3CB2 /* TransposeKernel.metal */; };
+		FCC15E04221E69E100DC3CB2 /* ConvAddPrelu.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DDB221E69E000DC3CB2 /* ConvAddPrelu.inc.metal */; };
+		FCC15E05221E69E100DC3CB2 /* BatchNormKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DDC221E69E000DC3CB2 /* BatchNormKernel.metal */; };
+		FCC15E06221E69E100DC3CB2 /* BoxCoder.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DDD221E69E000DC3CB2 /* BoxCoder.inc.metal */; };
+		FCC15E07221E69E100DC3CB2 /* FetchKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DDE221E69E000DC3CB2 /* FetchKernel.metal */; };
+		FCC15E08221E69E100DC3CB2 /* Split.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DDF221E69E000DC3CB2 /* Split.inc.metal */; };
+		FCC15E09221E69E100DC3CB2 /* ConcatKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DE0221E69E100DC3CB2 /* ConcatKernel.inc.metal */; };
+		FCC15E0A221E69E100DC3CB2 /* ElementwiseAddPreluKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DE1221E69E100DC3CB2 /* ElementwiseAddPreluKernel.inc.metal */; };
+		FCC15E0B221E69E100DC3CB2 /* FetchKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DE2221E69E100DC3CB2 /* FetchKernel.inc.metal */; };
+		FCC15E0C221E69E100DC3CB2 /* BufferToTexture.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DE3221E69E100DC3CB2 /* BufferToTexture.metal */; };
+		FCC15E0D221E69E100DC3CB2 /* ConvAddMetal.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCC15DE4221E69E100DC3CB2 /* ConvAddMetal.metal */; };
+/* End PBXBuildFile section */
+
+/* Begin PBXFileReference section */
+		33511F4FF7FE78679BE12DC0 /* Pods-paddle-mobile-metallib.release.xcconfig */ = {isa = PBXFileReference; includeInIndex = 1; lastKnownFileType = text.xcconfig; name = "Pods-paddle-mobile-metallib.release.xcconfig"; path = "../Pods/Target Support Files/Pods-paddle-mobile-metallib/Pods-paddle-mobile-metallib.release.xcconfig"; sourceTree = "<group>"; };
+		5D9D330A035906298947080B /* Pods_paddle_mobile_metallib.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; includeInIndex = 0; path = Pods_paddle_mobile_metallib.framework; sourceTree = BUILT_PRODUCTS_DIR; };
+		C6D31B9F9533810DBCA6B28D /* Pods-paddle-mobile-metallib.debug.xcconfig */ = {isa = PBXFileReference; includeInIndex = 1; lastKnownFileType = text.xcconfig; name = "Pods-paddle-mobile-metallib.debug.xcconfig"; path = "../Pods/Target Support Files/Pods-paddle-mobile-metallib/Pods-paddle-mobile-metallib.debug.xcconfig"; sourceTree = "<group>"; };
+		FCC15D60221E66DE00DC3CB2 /* paddle-mobile-metallib.metallib */ = {isa = PBXFileReference; explicitFileType = "archive.metal-library"; includeInIndex = 0; path = "paddle-mobile-metallib.metallib"; sourceTree = BUILT_PRODUCTS_DIR; };
+		FCC15DBC221E69DD00DC3CB2 /* ReluKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ReluKernel.metal; sourceTree = "<group>"; };
+		FCC15DBD221E69DD00DC3CB2 /* BoxCoder.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BoxCoder.metal; sourceTree = "<group>"; };
+		FCC15DBE221E69DD00DC3CB2 /* ConvAddBNReluKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConvAddBNReluKernel.metal; sourceTree = "<group>"; };
+		FCC15DBF221E69DD00DC3CB2 /* Split.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Split.metal; sourceTree = "<group>"; };
+		FCC15DC0221E69DD00DC3CB2 /* BilinearInterp.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BilinearInterp.metal; sourceTree = "<group>"; };
+		FCC15DC1221E69DD00DC3CB2 /* ElementwiseAddPreluKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ElementwiseAddPreluKernel.metal; sourceTree = "<group>"; };
+		FCC15DC2221E69DD00DC3CB2 /* NMSFetchResultKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = NMSFetchResultKernel.metal; sourceTree = "<group>"; };
+		FCC15DC3221E69DD00DC3CB2 /* Softmax.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Softmax.metal; sourceTree = "<group>"; };
+		FCC15DC4221E69DE00DC3CB2 /* PoolKernel.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = PoolKernel.inc.metal; sourceTree = "<group>"; };
+		FCC15DC5221E69DE00DC3CB2 /* ConvTransposeKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConvTransposeKernel.metal; sourceTree = "<group>"; };
+		FCC15DC6221E69DE00DC3CB2 /* Macro.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Macro.metal; sourceTree = "<group>"; };
+		FCC15DC7221E69DE00DC3CB2 /* PreluKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = PreluKernel.metal; sourceTree = "<group>"; };
+		FCC15DC8221E69DE00DC3CB2 /* BilinearInterp.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BilinearInterp.inc.metal; sourceTree = "<group>"; };
+		FCC15DC9221E69DE00DC3CB2 /* TransposeKernel.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = TransposeKernel.inc.metal; sourceTree = "<group>"; };
+		FCC15DCA221E69DE00DC3CB2 /* ConcatKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConcatKernel.metal; sourceTree = "<group>"; };
+		FCC15DCB221E69DE00DC3CB2 /* ResizeBilinear.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ResizeBilinear.metal; sourceTree = "<group>"; };
+		FCC15DCC221E69DE00DC3CB2 /* Common.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Common.metal; sourceTree = "<group>"; };
+		FCC15DCD221E69DE00DC3CB2 /* PoolKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = PoolKernel.metal; sourceTree = "<group>"; };
+		FCC15DCE221E69DE00DC3CB2 /* ReshapeKernel.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ReshapeKernel.inc.metal; sourceTree = "<group>"; };
+		FCC15DCF221E69DE00DC3CB2 /* ConvBNReluKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConvBNReluKernel.metal; sourceTree = "<group>"; };
+		FCC15DD0221E69DE00DC3CB2 /* Kernels.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Kernels.metal; sourceTree = "<group>"; };
+		FCC15DD1221E69DF00DC3CB2 /* Shape.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Shape.metal; sourceTree = "<group>"; };
+		FCC15DD2221E69DF00DC3CB2 /* Softmax.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Softmax.inc.metal; sourceTree = "<group>"; };
+		FCC15DD3221E69DF00DC3CB2 /* ConvAddPreluKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConvAddPreluKernel.metal; sourceTree = "<group>"; };
+		FCC15DD4221E69DF00DC3CB2 /* Elementwise.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Elementwise.metal; sourceTree = "<group>"; };
+		FCC15DD5221E69DF00DC3CB2 /* ReshapeKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ReshapeKernel.metal; sourceTree = "<group>"; };
+		FCC15DD6221E69DF00DC3CB2 /* Scale.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Scale.metal; sourceTree = "<group>"; };
+		FCC15DD7221E69DF00DC3CB2 /* ConvKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConvKernel.metal; sourceTree = "<group>"; };
+		FCC15DD8221E69DF00DC3CB2 /* PriorBoxKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = PriorBoxKernel.metal; sourceTree = "<group>"; };
+		FCC15DD9221E69E000DC3CB2 /* BatchNormRelu.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BatchNormRelu.metal; sourceTree = "<group>"; };
+		FCC15DDA221E69E000DC3CB2 /* TransposeKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = TransposeKernel.metal; sourceTree = "<group>"; };
+		FCC15DDB221E69E000DC3CB2 /* ConvAddPrelu.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConvAddPrelu.inc.metal; sourceTree = "<group>"; };
+		FCC15DDC221E69E000DC3CB2 /* BatchNormKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BatchNormKernel.metal; sourceTree = "<group>"; };
+		FCC15DDD221E69E000DC3CB2 /* BoxCoder.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BoxCoder.inc.metal; sourceTree = "<group>"; };
+		FCC15DDE221E69E000DC3CB2 /* FetchKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = FetchKernel.metal; sourceTree = "<group>"; };
+		FCC15DDF221E69E000DC3CB2 /* Split.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Split.inc.metal; sourceTree = "<group>"; };
+		FCC15DE0221E69E100DC3CB2 /* ConcatKernel.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConcatKernel.inc.metal; sourceTree = "<group>"; };
+		FCC15DE1221E69E100DC3CB2 /* ElementwiseAddPreluKernel.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ElementwiseAddPreluKernel.inc.metal; sourceTree = "<group>"; };
+		FCC15DE2221E69E100DC3CB2 /* FetchKernel.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = FetchKernel.inc.metal; sourceTree = "<group>"; };
+		FCC15DE3221E69E100DC3CB2 /* BufferToTexture.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BufferToTexture.metal; sourceTree = "<group>"; };
+		FCC15DE4221E69E100DC3CB2 /* ConvAddMetal.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConvAddMetal.metal; sourceTree = "<group>"; };
+/* End PBXFileReference section */
+
+/* Begin PBXFrameworksBuildPhase section */
+		3262940821E130A79BEC3880 /* Frameworks */ = {
+			isa = PBXFrameworksBuildPhase;
+			buildActionMask = 2147483647;
+			files = (
+				5CCC0CF6759710BAFE999DB7 /* Pods_paddle_mobile_metallib.framework in Frameworks */,
+			);
+			runOnlyForDeploymentPostprocessing = 0;
+		};
+/* End PBXFrameworksBuildPhase section */
+
+/* Begin PBXGroup section */
+		2EC9E2608C2591494F8F23F0 /* Frameworks */ = {
+			isa = PBXGroup;
+			children = (
+				5D9D330A035906298947080B /* Pods_paddle_mobile_metallib.framework */,
+			);
+			name = Frameworks;
+			sourceTree = "<group>";
+		};
+		755C26B34D5114CE1B98D3DC /* Pods */ = {
+			isa = PBXGroup;
+			children = (
+				C6D31B9F9533810DBCA6B28D /* Pods-paddle-mobile-metallib.debug.xcconfig */,
+				33511F4FF7FE78679BE12DC0 /* Pods-paddle-mobile-metallib.release.xcconfig */,
+			);
+			name = Pods;
+			sourceTree = "<group>";
+		};
+		FCC15D59221E66DE00DC3CB2 = {
+			isa = PBXGroup;
+			children = (
+				FCC15D62221E66DE00DC3CB2 /* paddle-mobile-metallib */,
+				FCC15D61221E66DE00DC3CB2 /* Products */,
+				755C26B34D5114CE1B98D3DC /* Pods */,
+				2EC9E2608C2591494F8F23F0 /* Frameworks */,
+			);
+			sourceTree = "<group>";
+		};
+		FCC15D61221E66DE00DC3CB2 /* Products */ = {
+			isa = PBXGroup;
+			children = (
+				FCC15D60221E66DE00DC3CB2 /* paddle-mobile-metallib.metallib */,
+			);
+			name = Products;
+			sourceTree = "<group>";
+		};
+		FCC15D62221E66DE00DC3CB2 /* paddle-mobile-metallib */ = {
+			isa = PBXGroup;
+			children = (
+				FCC15DDC221E69E000DC3CB2 /* BatchNormKernel.metal */,
+				FCC15DD9221E69E000DC3CB2 /* BatchNormRelu.metal */,
+				FCC15DC8221E69DE00DC3CB2 /* BilinearInterp.inc.metal */,
+				FCC15DC0221E69DD00DC3CB2 /* BilinearInterp.metal */,
+				FCC15DDD221E69E000DC3CB2 /* BoxCoder.inc.metal */,
+				FCC15DBD221E69DD00DC3CB2 /* BoxCoder.metal */,
+				FCC15DE3221E69E100DC3CB2 /* BufferToTexture.metal */,
+				FCC15DCC221E69DE00DC3CB2 /* Common.metal */,
+				FCC15DE0221E69E100DC3CB2 /* ConcatKernel.inc.metal */,
+				FCC15DCA221E69DE00DC3CB2 /* ConcatKernel.metal */,
+				FCC15DBE221E69DD00DC3CB2 /* ConvAddBNReluKernel.metal */,
+				FCC15DE4221E69E100DC3CB2 /* ConvAddMetal.metal */,
+				FCC15DDB221E69E000DC3CB2 /* ConvAddPrelu.inc.metal */,
+				FCC15DD3221E69DF00DC3CB2 /* ConvAddPreluKernel.metal */,
+				FCC15DCF221E69DE00DC3CB2 /* ConvBNReluKernel.metal */,
+				FCC15DD7221E69DF00DC3CB2 /* ConvKernel.metal */,
+				FCC15DC5221E69DE00DC3CB2 /* ConvTransposeKernel.metal */,
+				FCC15DD4221E69DF00DC3CB2 /* Elementwise.metal */,
+				FCC15DE1221E69E100DC3CB2 /* ElementwiseAddPreluKernel.inc.metal */,
+				FCC15DC1221E69DD00DC3CB2 /* ElementwiseAddPreluKernel.metal */,
+				FCC15DE2221E69E100DC3CB2 /* FetchKernel.inc.metal */,
+				FCC15DDE221E69E000DC3CB2 /* FetchKernel.metal */,
+				FCC15DD0221E69DE00DC3CB2 /* Kernels.metal */,
+				FCC15DC6221E69DE00DC3CB2 /* Macro.metal */,
+				FCC15DC2221E69DD00DC3CB2 /* NMSFetchResultKernel.metal */,
+				FCC15DC4221E69DE00DC3CB2 /* PoolKernel.inc.metal */,
+				FCC15DCD221E69DE00DC3CB2 /* PoolKernel.metal */,
+				FCC15DC7221E69DE00DC3CB2 /* PreluKernel.metal */,
+				FCC15DD8221E69DF00DC3CB2 /* PriorBoxKernel.metal */,
+				FCC15DBC221E69DD00DC3CB2 /* ReluKernel.metal */,
+				FCC15DCE221E69DE00DC3CB2 /* ReshapeKernel.inc.metal */,
+				FCC15DD5221E69DF00DC3CB2 /* ReshapeKernel.metal */,
+				FCC15DCB221E69DE00DC3CB2 /* ResizeBilinear.metal */,
+				FCC15DD6221E69DF00DC3CB2 /* Scale.metal */,
+				FCC15DD1221E69DF00DC3CB2 /* Shape.metal */,
+				FCC15DD2221E69DF00DC3CB2 /* Softmax.inc.metal */,
+				FCC15DC3221E69DD00DC3CB2 /* Softmax.metal */,
+				FCC15DDF221E69E000DC3CB2 /* Split.inc.metal */,
+				FCC15DBF221E69DD00DC3CB2 /* Split.metal */,
+				FCC15DC9221E69DE00DC3CB2 /* TransposeKernel.inc.metal */,
+				FCC15DDA221E69E000DC3CB2 /* TransposeKernel.metal */,
+			);
+			path = "paddle-mobile-metallib";
+			sourceTree = "<group>";
+		};
+/* End PBXGroup section */
+
+/* Begin PBXNativeTarget section */
+		FCC15D5F221E66DE00DC3CB2 /* paddle-mobile-metallib */ = {
+			isa = PBXNativeTarget;
+			buildConfigurationList = FCC15D67221E66DE00DC3CB2 /* Build configuration list for PBXNativeTarget "paddle-mobile-metallib" */;
+			buildPhases = (
+				DD854B7EC9A77A557887A67F /* [CP] Check Pods Manifest.lock */,
+				FCC15D5E221E66DE00DC3CB2 /* Sources */,
+				3262940821E130A79BEC3880 /* Frameworks */,
+			);
+			buildRules = (
+			);
+			dependencies = (
+			);
+			name = "paddle-mobile-metallib";
+			productName = "paddle-mobile-metallib";
+			productReference = FCC15D60221E66DE00DC3CB2 /* paddle-mobile-metallib.metallib */;
+			productType = "com.apple.product-type.metal-library";
+		};
+/* End PBXNativeTarget section */
+
+/* Begin PBXProject section */
+		FCC15D5A221E66DE00DC3CB2 /* Project object */ = {
+			isa = PBXProject;
+			attributes = {
+				LastUpgradeCheck = 1010;
+				ORGANIZATIONNAME = Ray;
+				TargetAttributes = {
+					FCC15D5F221E66DE00DC3CB2 = {
+						CreatedOnToolsVersion = 10.1;
+					};
+				};
+			};
+			buildConfigurationList = FCC15D5D221E66DE00DC3CB2 /* Build configuration list for PBXProject "paddle-mobile-metallib" */;
+			compatibilityVersion = "Xcode 9.3";
+			developmentRegion = en;
+			hasScannedForEncodings = 0;
+			knownRegions = (
+				en,
+			);
+			mainGroup = FCC15D59221E66DE00DC3CB2;
+			productRefGroup = FCC15D61221E66DE00DC3CB2 /* Products */;
+			projectDirPath = "";
+			projectRoot = "";
+			targets = (
+				FCC15D5F221E66DE00DC3CB2 /* paddle-mobile-metallib */,
+			);
+		};
+/* End PBXProject section */
+
+/* Begin PBXShellScriptBuildPhase section */
+		DD854B7EC9A77A557887A67F /* [CP] Check Pods Manifest.lock */ = {
+			isa = PBXShellScriptBuildPhase;
+			buildActionMask = 2147483647;
+			files = (
+			);
+			inputFileListPaths = (
+			);
+			inputPaths = (
+				"${PODS_PODFILE_DIR_PATH}/Podfile.lock",
+				"${PODS_ROOT}/Manifest.lock",
+			);
+			name = "[CP] Check Pods Manifest.lock";
+			outputFileListPaths = (
+			);
+			outputPaths = (
+				"$(DERIVED_FILE_DIR)/Pods-paddle-mobile-metallib-checkManifestLockResult.txt",
+			);
+			runOnlyForDeploymentPostprocessing = 0;
+			shellPath = /bin/sh;
+			shellScript = "diff \"${PODS_PODFILE_DIR_PATH}/Podfile.lock\" \"${PODS_ROOT}/Manifest.lock\" > /dev/null\nif [ $? != 0 ] ; then\n    # print error to STDERR\n    echo \"error: The sandbox is not in sync with the Podfile.lock. Run 'pod install' or update your CocoaPods installation.\" >&2\n    exit 1\nfi\n# This output is used by Xcode 'outputs' to avoid re-running this script phase.\necho \"SUCCESS\" > \"${SCRIPT_OUTPUT_FILE_0}\"\n";
+			showEnvVarsInLog = 0;
+		};
+/* End PBXShellScriptBuildPhase section */
+
+/* Begin PBXSourcesBuildPhase section */
+		FCC15D5E221E66DE00DC3CB2 /* Sources */ = {
+			isa = PBXSourcesBuildPhase;
+			buildActionMask = 2147483647;
+			files = (
+				FCC15DF5221E69E100DC3CB2 /* Common.metal in Sources */,
+				FCC15DFF221E69E100DC3CB2 /* Scale.metal in Sources */,
+				FCC15DEC221E69E100DC3CB2 /* Softmax.metal in Sources */,
+				FCC15DE8221E69E100DC3CB2 /* Split.metal in Sources */,
+				FCC15DF2221E69E100DC3CB2 /* TransposeKernel.inc.metal in Sources */,
+				FCC15DE7221E69E100DC3CB2 /* ConvAddBNReluKernel.metal in Sources */,
+				FCC15E04221E69E100DC3CB2 /* ConvAddPrelu.inc.metal in Sources */,
+				FCC15DF9221E69E100DC3CB2 /* Kernels.metal in Sources */,
+				FCC15DF0221E69E100DC3CB2 /* PreluKernel.metal in Sources */,
+				FCC15DEB221E69E100DC3CB2 /* NMSFetchResultKernel.metal in Sources */,
+				FCC15DE9221E69E100DC3CB2 /* BilinearInterp.metal in Sources */,
+				FCC15DFA221E69E100DC3CB2 /* Shape.metal in Sources */,
+				FCC15E0C221E69E100DC3CB2 /* BufferToTexture.metal in Sources */,
+				FCC15E0A221E69E100DC3CB2 /* ElementwiseAddPreluKernel.inc.metal in Sources */,
+				FCC15DFB221E69E100DC3CB2 /* Softmax.inc.metal in Sources */,
+				FCC15E03221E69E100DC3CB2 /* TransposeKernel.metal in Sources */,
+				FCC15DFE221E69E100DC3CB2 /* ReshapeKernel.metal in Sources */,
+				FCC15E0D221E69E100DC3CB2 /* ConvAddMetal.metal in Sources */,
+				FCC15DF7221E69E100DC3CB2 /* ReshapeKernel.inc.metal in Sources */,
+				FCC15DE5221E69E100DC3CB2 /* ReluKernel.metal in Sources */,
+				FCC15DEF221E69E100DC3CB2 /* Macro.metal in Sources */,
+				FCC15E02221E69E100DC3CB2 /* BatchNormRelu.metal in Sources */,
+				FCC15E00221E69E100DC3CB2 /* ConvKernel.metal in Sources */,
+				FCC15E01221E69E100DC3CB2 /* PriorBoxKernel.metal in Sources */,
+				FCC15DEA221E69E100DC3CB2 /* ElementwiseAddPreluKernel.metal in Sources */,
+				FCC15DED221E69E100DC3CB2 /* PoolKernel.inc.metal in Sources */,
+				FCC15E07221E69E100DC3CB2 /* FetchKernel.metal in Sources */,
+				FCC15E0B221E69E100DC3CB2 /* FetchKernel.inc.metal in Sources */,
+				FCC15DEE221E69E100DC3CB2 /* ConvTransposeKernel.metal in Sources */,
+				FCC15DFC221E69E100DC3CB2 /* ConvAddPreluKernel.metal in Sources */,
+				FCC15E06221E69E100DC3CB2 /* BoxCoder.inc.metal in Sources */,
+				FCC15DF1221E69E100DC3CB2 /* BilinearInterp.inc.metal in Sources */,
+				FCC15E08221E69E100DC3CB2 /* Split.inc.metal in Sources */,
+				FCC15DF4221E69E100DC3CB2 /* ResizeBilinear.metal in Sources */,
+				FCC15E05221E69E100DC3CB2 /* BatchNormKernel.metal in Sources */,
+				FCC15DE6221E69E100DC3CB2 /* BoxCoder.metal in Sources */,
+				FCC15DF6221E69E100DC3CB2 /* PoolKernel.metal in Sources */,
+				FCC15E09221E69E100DC3CB2 /* ConcatKernel.inc.metal in Sources */,
+				FCC15DFD221E69E100DC3CB2 /* Elementwise.metal in Sources */,
+				FCC15DF8221E69E100DC3CB2 /* ConvBNReluKernel.metal in Sources */,
+				FCC15DF3221E69E100DC3CB2 /* ConcatKernel.metal in Sources */,
+			);
+			runOnlyForDeploymentPostprocessing = 0;
+		};
+/* End PBXSourcesBuildPhase section */
+
+/* Begin XCBuildConfiguration section */
+		FCC15D65221E66DE00DC3CB2 /* Debug */ = {
+			isa = XCBuildConfiguration;
+			buildSettings = {
+				ALWAYS_SEARCH_USER_PATHS = NO;
+				IPHONEOS_DEPLOYMENT_TARGET = 12.1;
+				MTL_ENABLE_DEBUG_INFO = INCLUDE_SOURCE;
+				MTL_FAST_MATH = YES;
+				SDKROOT = iphoneos;
+			};
+			name = Debug;
+		};
+		FCC15D66221E66DE00DC3CB2 /* Release */ = {
+			isa = XCBuildConfiguration;
+			buildSettings = {
+				ALWAYS_SEARCH_USER_PATHS = NO;
+				IPHONEOS_DEPLOYMENT_TARGET = 12.1;
+				MTL_ENABLE_DEBUG_INFO = NO;
+				MTL_FAST_MATH = YES;
+				SDKROOT = iphoneos;
+			};
+			name = Release;
+		};
+		FCC15D68221E66DE00DC3CB2 /* Debug */ = {
+			isa = XCBuildConfiguration;
+			baseConfigurationReference = C6D31B9F9533810DBCA6B28D /* Pods-paddle-mobile-metallib.debug.xcconfig */;
+			buildSettings = {
+				CODE_SIGN_STYLE = Automatic;
+				MTL_ENABLE_DEBUG_INFO = NO;
+				PRODUCT_NAME = "$(TARGET_NAME)";
+			};
+			name = Debug;
+		};
+		FCC15D69221E66DE00DC3CB2 /* Release */ = {
+			isa = XCBuildConfiguration;
+			baseConfigurationReference = 33511F4FF7FE78679BE12DC0 /* Pods-paddle-mobile-metallib.release.xcconfig */;
+			buildSettings = {
+				CODE_SIGN_STYLE = Automatic;
+				MTL_ENABLE_DEBUG_INFO = NO;
+				PRODUCT_NAME = "$(TARGET_NAME)";
+			};
+			name = Release;
+		};
+/* End XCBuildConfiguration section */
+
+/* Begin XCConfigurationList section */
+		FCC15D5D221E66DE00DC3CB2 /* Build configuration list for PBXProject "paddle-mobile-metallib" */ = {
+			isa = XCConfigurationList;
+			buildConfigurations = (
+				FCC15D65221E66DE00DC3CB2 /* Debug */,
+				FCC15D66221E66DE00DC3CB2 /* Release */,
+			);
+			defaultConfigurationIsVisible = 0;
+			defaultConfigurationName = Release;
+		};
+		FCC15D67221E66DE00DC3CB2 /* Build configuration list for PBXNativeTarget "paddle-mobile-metallib" */ = {
+			isa = XCConfigurationList;
+			buildConfigurations = (
+				FCC15D68221E66DE00DC3CB2 /* Debug */,
+				FCC15D69221E66DE00DC3CB2 /* Release */,
+			);
+			defaultConfigurationIsVisible = 0;
+			defaultConfigurationName = Release;
+		};
+/* End XCConfigurationList section */
+	};
+	rootObject = FCC15D5A221E66DE00DC3CB2 /* Project object */;
+}

metal/paddle-mobile-metallib/paddle-mobile-metallib.xcodeproj/project.xcworkspace/contents.xcworkspacedata

+<?xml version="1.0" encoding="UTF-8"?>
+<Workspace
+   version = "1.0">
+   <FileRef
+      location = "self:paddle-mobile-metallib.xcodeproj">
+   </FileRef>
+</Workspace>

metal/paddle-mobile-metallib/paddle-mobile-metallib.xcodeproj/project.xcworkspace/xcshareddata/IDEWorkspaceChecks.plist

+<?xml version="1.0" encoding="UTF-8"?>
+<!DOCTYPE plist PUBLIC "-//Apple//DTD PLIST 1.0//EN" "http://www.apple.com/DTDs/PropertyList-1.0.dtd">
+<plist version="1.0">
+<dict>
+	<key>IDEDidComputeMac32BitWarning</key>
+	<true/>
+</dict>
+</plist>

metal/paddle-mobile-metallib/paddle-mobile-metallib.xcodeproj/xcshareddata/xcschemes/paddle-mobile-metallib.xcscheme

+<?xml version="1.0" encoding="UTF-8"?>
+<Scheme
+   LastUpgradeVersion = "1010"
+   version = "1.3">
+   <BuildAction
+      parallelizeBuildables = "YES"
+      buildImplicitDependencies = "YES">
+      <BuildActionEntries>
+         <BuildActionEntry
+            buildForTesting = "YES"
+            buildForRunning = "YES"
+            buildForProfiling = "YES"
+            buildForArchiving = "YES"
+            buildForAnalyzing = "YES">
+            <BuildableReference
+               BuildableIdentifier = "primary"
+               BlueprintIdentifier = "FCC15D5F221E66DE00DC3CB2"
+               BuildableName = "paddle-mobile-metallib.metallib"
+               BlueprintName = "paddle-mobile-metallib"
+               ReferencedContainer = "container:paddle-mobile-metallib.xcodeproj">
+            </BuildableReference>
+         </BuildActionEntry>
+      </BuildActionEntries>
+   </BuildAction>
+   <TestAction
+      buildConfiguration = "Debug"
+      selectedDebuggerIdentifier = "Xcode.DebuggerFoundation.Debugger.LLDB"
+      selectedLauncherIdentifier = "Xcode.DebuggerFoundation.Launcher.LLDB"
+      shouldUseLaunchSchemeArgsEnv = "YES">
+      <Testables>
+      </Testables>
+      <AdditionalOptions>
+      </AdditionalOptions>
+   </TestAction>
+   <LaunchAction
+      buildConfiguration = "Release"
+      selectedDebuggerIdentifier = "Xcode.DebuggerFoundation.Debugger.LLDB"
+      selectedLauncherIdentifier = "Xcode.DebuggerFoundation.Launcher.LLDB"
+      launchStyle = "0"
+      useCustomWorkingDirectory = "NO"
+      ignoresPersistentStateOnLaunch = "NO"
+      debugDocumentVersioning = "YES"
+      debugServiceExtension = "internal"
+      allowLocationSimulation = "YES">
+      <MacroExpansion>
+         <BuildableReference
+            BuildableIdentifier = "primary"
+            BlueprintIdentifier = "FCC15D5F221E66DE00DC3CB2"
+            BuildableName = "paddle-mobile-metallib.metallib"
+            BlueprintName = "paddle-mobile-metallib"
+            ReferencedContainer = "container:paddle-mobile-metallib.xcodeproj">
+         </BuildableReference>
+      </MacroExpansion>
+      <AdditionalOptions>
+      </AdditionalOptions>
+   </LaunchAction>
+   <ProfileAction
+      buildConfiguration = "Release"
+      shouldUseLaunchSchemeArgsEnv = "YES"
+      savedToolIdentifier = ""
+      useCustomWorkingDirectory = "NO"
+      debugDocumentVersioning = "YES">
+      <MacroExpansion>
+         <BuildableReference
+            BuildableIdentifier = "primary"
+            BlueprintIdentifier = "FCC15D5F221E66DE00DC3CB2"
+            BuildableName = "paddle-mobile-metallib.metallib"
+            BlueprintName = "paddle-mobile-metallib"
+            ReferencedContainer = "container:paddle-mobile-metallib.xcodeproj">
+         </BuildableReference>
+      </MacroExpansion>
+   </ProfileAction>
+   <AnalyzeAction
+      buildConfiguration = "Debug">
+   </AnalyzeAction>
+   <ArchiveAction
+      buildConfiguration = "Release"
+      revealArchiveInOrganizer = "YES">
+   </ArchiveAction>
+</Scheme>

metal/paddle-mobile-metallib/paddle-mobile-metallib/BatchNormRelu.metal

+//
+//  BatchNormRelu.metal
+//  paddle-mobile
+//
+
+#include <metal_stdlib>
+using namespace metal;
+
+struct MetalConvParam {
+    short offsetX;
+    short offsetY;
+    short offsetZ;
+    ushort strideX;
+    ushort strideY;
+};
+
+kernel void batch_norm_relu_3x3(texture2d_array<float, access::sample> inTexture [[texture(0)]],
+                                         texture2d_array<float, access::write> outTexture [[texture(1)]],
+                                         const device float4 *new_scale [[buffer(0)]],
+                                         const device float4 *new_biase [[buffer(1)]],
+                                         uint3 gid [[thread_position_in_grid]]) {
+    
+    if (gid.x >= outTexture.get_width() ||
+        gid.y >= outTexture.get_height() ||
+        gid.z >= outTexture.get_array_size()) {
+        return;
+    }
+    
+    float4 input;
+    float4 output;
+    constexpr sampler sample(coord::pixel, filter::nearest, address::clamp_to_zero);
+    input = inTexture.sample(sample, gid.x, gid.y, gid.z);
+    output = fmax(input * new_scale[gid.z] + new_biase[gid.z], 0.0);
+    outTexture.write(output, gid.xy, gid.z);
+
+}

metal/paddle-mobile-metallib/paddle-mobile-metallib/ResizeBilinear.metal

+/* 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>
+using namespace metal;
+
+struct resize_bilinear_param {
+//  int32_t out_h;
+//  int32_t out_w;
+  float ratio_h;
+  float ratio_w;
+};
+
+kernel void resize_bilinear(texture2d_array<float, access::read> input [[texture(0)]],
+                     texture2d_array<float, access::write> output [[texture(2)]],
+                     constant resize_bilinear_param & pm [[buffer(0)]],
+                     uint3 gid [[thread_position_in_grid]]) {
+  float4 r;
+  if ((input.get_width() == output.get_width()) && (input.get_height() == output.get_height())) {
+    r = input.read(gid.xy, gid.z);
+  } else {
+    float w = gid.x * pm.ratio_w;
+    float h = gid.y * pm.ratio_h;
+    uint w0 = w, h0 = h;
+    uint w1 = w0 + 1, h1 = h0 + 1;
+    float w1lambda = w - w0, h1lambda = h - h0;
+    float w2lambda = 1.0 - w1lambda, h2lambda = 1.0 - h1lambda;
+    if (w1 >= input.get_width()) w1 = w0;
+    if (h1 >= input.get_height()) h1 = h0;
+    float4 r0 = input.read(uint2(w0, h0), gid.z);
+    float4 r1 = input.read(uint2(w1, h0), gid.z);
+    float4 r2 = input.read(uint2(w0, h1), gid.z);
+    float4 r3 = input.read(uint2(w1, h1), gid.z);
+    r = h2lambda * (w2lambda * r0 + w1lambda * r1) + h1lambda * (w2lambda * r2 + w1lambda * r3);
+  }
+  output.write(r, gid.xy, gid.z);
+}
+
+kernel void resize_bilinear_half(texture2d_array<half, access::read> input [[texture(0)]],
+                            texture2d_array<half, access::write> output [[texture(2)]],
+                            constant resize_bilinear_param & pm [[buffer(0)]],
+                            uint3 gid [[thread_position_in_grid]]) {
+  
+  half4 r;
+  if ((input.get_width() == output.get_width()) && (input.get_height() == output.get_height())) {
+    r = input.read(gid.xy, gid.z);
+  } else {
+    half w = gid.x * pm.ratio_w;
+    half h = gid.y * pm.ratio_h;
+    uint w0 = w, h0 = h;
+    uint w1 = w0 + 1, h1 = h0 + 1;
+    half w1lambda = w - w0, h1lambda = h - h0;
+    half w2lambda = 1.0 - w1lambda, h2lambda = 1.0 - h1lambda;
+    if (w1 >= input.get_width()) w1 = w0;
+    if (h1 >= input.get_height()) h1 = h0;
+    half4 r0 = input.read(uint2(w0, h0), gid.z);
+    half4 r1 = input.read(uint2(w1, h0), gid.z);
+    half4 r2 = input.read(uint2(w0, h1), gid.z);
+    half4 r3 = input.read(uint2(w1, h1), gid.z);
+    r = h2lambda * (w2lambda * r0 + w1lambda * r1) + h1lambda * (w2lambda * r2 + w1lambda * r3);
+  }
+  output.write(r, gid.xy, gid.z);
+  output.write(r, gid.xy, gid.z);
+}

metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj

@@ -11,31 +11,16 @@
 		456BB7B521F5B356001474E2 /* Framework.pbobjc.h in Headers */ = {isa = PBXBuildFile; fileRef = 456BB7B321F5B356001474E2 /* Framework.pbobjc.h */; settings = {ATTRIBUTES = (Public, ); }; };
 		4AA1EA862146625E00D0F791 /* BilinearInterpOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA852146625E00D0F791 /* BilinearInterpOp.swift */; };
 		4AA1EA88214662BD00D0F791 /* BilinearInterpKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA87214662BD00D0F791 /* BilinearInterpKernel.swift */; };
-		4AA1EA8A2146631C00D0F791 /* BilinearInterp.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA892146631C00D0F791 /* BilinearInterp.metal */; };
 		4AA1EA8C2146640900D0F791 /* SplitOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA8B2146640900D0F791 /* SplitOp.swift */; };
 		4AA1EA8E2146647F00D0F791 /* SplitKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA8D2146647F00D0F791 /* SplitKernel.swift */; };
-		4AA1EA90214664CD00D0F791 /* Split.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA8F214664CD00D0F791 /* Split.metal */; };
 		4AA1EA92214665D700D0F791 /* ShapeOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA91214665D700D0F791 /* ShapeOp.swift */; };
 		4AA1EA942146661500D0F791 /* ShapeKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA932146661500D0F791 /* ShapeKernel.swift */; };
 		4AA1EA982146666500D0F791 /* FlattenOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA972146666500D0F791 /* FlattenOp.swift */; };
-		4AA1EAA02148DEEE00D0F791 /* ReshapeKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA9F2148DEEE00D0F791 /* ReshapeKernel.inc.metal */; };
 		4AA1EAA2214912CD00D0F791 /* FlattenKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EAA1214912CC00D0F791 /* FlattenKernel.swift */; };
-		4AA1EAA4214A295C00D0F791 /* Split.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EAA3214A295C00D0F791 /* Split.inc.metal */; };
-		4AA1EAA6214B5F6800D0F791 /* Shape.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EAA5214B5F6800D0F791 /* Shape.metal */; };
-		4AA1EAA8214B7AFB00D0F791 /* BilinearInterp.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EAA7214B7AFB00D0F791 /* BilinearInterp.inc.metal */; };
-		4AA1EAAA214F53D800D0F791 /* BoxCoder.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EAA9214F53D800D0F791 /* BoxCoder.inc.metal */; };
-		4AA1EAAC214F55C800D0F791 /* Softmax.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EAAB214F55C800D0F791 /* Softmax.inc.metal */; };
-		4AA1EAAE214F5FD900D0F791 /* TransposeKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EAAD214F5FD900D0F791 /* TransposeKernel.inc.metal */; };
-		4AF928772133F1DB005B6C3A /* BoxCoder.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AF928762133F1DB005B6C3A /* BoxCoder.metal */; };
-		4AF9287921341661005B6C3A /* Softmax.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AF9287821341661005B6C3A /* Softmax.metal */; };
-		4AF928822135673D005B6C3A /* ConcatKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AF928812135673D005B6C3A /* ConcatKernel.metal */; };
-		4AF9288421357BE3005B6C3A /* Elementwise.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AF9288321357BE3005B6C3A /* Elementwise.metal */; };
 		C28FE02F21BA68C00054EFAC /* Metal.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = C28FE02C21BA68C00054EFAC /* Metal.framework */; };
 		C28FE03021BA68C00054EFAC /* MetalPerformanceShaders.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = C28FE02D21BA68C00054EFAC /* MetalPerformanceShaders.framework */; };
 		C28FE03121BA68C00054EFAC /* MetalKit.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = C28FE02E21BA68C00054EFAC /* MetalKit.framework */; };
 		D3831F70E7E0B565B9AC22DA /* Pods_paddle_mobile.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = DD2E06330A1E7129C918DB46 /* Pods_paddle_mobile.framework */; };
-		FC0226562138F33800F395E2 /* TransposeKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC0226552138F33800F395E2 /* TransposeKernel.metal */; };
-		FC0226582138F38D00F395E2 /* PoolKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC0226572138F38D00F395E2 /* PoolKernel.metal */; };
 		FC039B6F20E11C3C0081E9F8 /* paddle_mobile.h in Headers */ = {isa = PBXBuildFile; fileRef = FC039B6D20E11C3C0081E9F8 /* paddle_mobile.h */; settings = {ATTRIBUTES = (Public, ); }; };
 		FC039B9720E11C9A0081E9F8 /* Extensions.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC039B9420E11C9A0081E9F8 /* Extensions.swift */; };
 		FC039B9820E11C9A0081E9F8 /* Errors.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC039B9520E11C9A0081E9F8 /* Errors.swift */; };
@@ -61,40 +46,25 @@
 		FC0E2DBC20EE45FE009C1FAC /* ConvKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC0E2DBB20EE45FE009C1FAC /* ConvKernel.swift */; };
 		FC0E2DBE20EE460D009C1FAC /* BatchNormKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC0E2DBD20EE460D009C1FAC /* BatchNormKernel.swift */; };
 		FC0E2DC020EE461F009C1FAC /* ElementwiseAddKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */; };
-		FC1B16B320EC9A4F00678B91 /* Kernels.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC1B16B220EC9A4F00678B91 /* Kernels.metal */; };
 		FC1CF3F721D4B4C400F7392E /* Runner.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC1CF3F621D4B4C400F7392E /* Runner.swift */; };
 		FC2BFCC221DF2F9100C262B2 /* GlobalConfig.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC2BFCC121DF2F9100C262B2 /* GlobalConfig.swift */; };
 		FC2BFD4621DF685F00C262B2 /* Scale.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC2BFD4521DF685F00C262B2 /* Scale.swift */; };
 		FC2BFD4A21DF81DE00C262B2 /* Kernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC2BFD4921DF81DE00C262B2 /* Kernel.swift */; };
 		FC2BFD4E21DF820B00C262B2 /* ConvAddBatchNormReluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC2BFD4D21DF820A00C262B2 /* ConvAddBatchNormReluOp.swift */; };
-		FC2BFD5121DF8E0400C262B2 /* Scale.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC2BFD5021DF8E0400C262B2 /* Scale.metal */; };
 		FC3602CC2108819F00FACB58 /* PaddleMobileUnitTest.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */; };
-		FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74820F0B954007C0C6D /* ConvKernel.metal */; };
 		FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */; };
 		FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */; };
 		FC60DB8920E9AAA500FF203F /* MetalExtension.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC60DB8820E9AAA500FF203F /* MetalExtension.swift */; };
 		FC803BBF214CB65A0094B8E5 /* ConvAddPreluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC803BBE214CB65A0094B8E5 /* ConvAddPreluOp.swift */; };
 		FC803BC1214CB77A0094B8E5 /* ConvAddPreluKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC803BC0214CB77A0094B8E5 /* ConvAddPreluKernel.swift */; };
-		FC803BC3214CB79C0094B8E5 /* ConvAddPreluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC803BC2214CB79C0094B8E5 /* ConvAddPreluKernel.metal */; };
-		FC803BC5214CB8F00094B8E5 /* ConvAddPrelu.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC803BC4214CB8F00094B8E5 /* ConvAddPrelu.inc.metal */; };
-		FC803BC7214CBA820094B8E5 /* Macro.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC803BC6214CBA820094B8E5 /* Macro.metal */; };
-		FC803BC9214CFC8D0094B8E5 /* FetchKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC803BC8214CFC8D0094B8E5 /* FetchKernel.metal */; };
 		FC82735920E3C04200BE430A /* OpCreator.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC82735820E3C04200BE430A /* OpCreator.swift */; };
 		FC9797C921D6101D00F2FD90 /* ResizeBilinearOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9797C821D6101D00F2FD90 /* ResizeBilinearOp.swift */; };
 		FC9797CB21D6102D00F2FD90 /* ResizeBilinearKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9797CA21D6102D00F2FD90 /* ResizeBilinearKernel.swift */; };
-		FC9C2A0D21D3D185005856C6 /* FetchKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC9C2A0C21D3D185005856C6 /* FetchKernel.inc.metal */; };
 		FC9D037920E229E4000F735A /* OpParam.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037820E229E4000F735A /* OpParam.swift */; };
 		FC9D038020E22FBB000F735A /* FeedOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037F20E22FBB000F735A /* FeedOp.swift */; };
 		FC9D038220E2312E000F735A /* FetchOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D038120E2312E000F735A /* FetchOp.swift */; };
 		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 */; };
-		FCA67B1721364EF000BD58AA /* ConvTransposeKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCA67B1621364EF000BD58AA /* ConvTransposeKernel.metal */; };
-		FCA67CD52138272900BD58AA /* ConvAddMetal.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCA67CD42138272900BD58AA /* ConvAddMetal.metal */; };
-		FCA67CD7213827AC00BD58AA /* ConvAddBNReluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCA67CD6213827AC00BD58AA /* ConvAddBNReluKernel.metal */; };
-		FCA67CD92138287B00BD58AA /* ConvBNReluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCA67CD82138287B00BD58AA /* ConvBNReluKernel.metal */; };
 		FCB40E5921E0DCAB0075EC91 /* FetchKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCB40E5821E0DCAB0075EC91 /* FetchKernel.swift */; };
-		FCB91DC221FEEE990051C6B2 /* BufferToTexture.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCB91DC121FEEE990051C6B2 /* BufferToTexture.metal */; };
 		FCBCCC572122F41300D94F7E /* DwConvBNReluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC562122F41300D94F7E /* DwConvBNReluOp.swift */; };
 		FCBCCC592122F42700D94F7E /* ConvBNReluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC582122F42700D94F7E /* ConvBNReluOp.swift */; };
 		FCBCCC5B2122F66F00D94F7E /* ConvBNReluKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC5A2122F66F00D94F7E /* ConvBNReluKernel.swift */; };
@@ -109,7 +79,6 @@
 		FCBCCC6D2123073A00D94F7E /* BoxcoderKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC6C2123073A00D94F7E /* BoxcoderKernel.swift */; };
 		FCBCCC6F2123097100D94F7E /* MulticlassNMSOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC6E2123097100D94F7E /* MulticlassNMSOp.swift */; };
 		FCBCCC71212309A700D94F7E /* MulticlassNMSKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC70212309A700D94F7E /* MulticlassNMSKernel.swift */; };
-		FCCED5E121D71FC000BE8D5F /* PoolKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCCED5E021D71FC000BE8D5F /* PoolKernel.inc.metal */; };
 		FCD04E6620F314C50007374F /* PoolOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCD04E6520F314C50007374F /* PoolOp.swift */; };
 		FCD04E6820F315020007374F /* PoolKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCD04E6720F315020007374F /* PoolKernel.swift */; };
 		FCD04E6A20F319EC0007374F /* SoftmaxOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCD04E6920F319EC0007374F /* SoftmaxOp.swift */; };
@@ -121,19 +90,12 @@
 		FCDC0FEB21099A1D00DC9EFB /* Tools.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCDC0FEA21099A1D00DC9EFB /* Tools.swift */; };
 		FCDDC6C6212F9FB800E5EF74 /* PreluKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCDDC6C5212F9FB800E5EF74 /* PreluKernel.swift */; };
 		FCDDC6C8212FA3CA00E5EF74 /* ConvTransposeKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCDDC6C7212FA3CA00E5EF74 /* ConvTransposeKernel.swift */; };
-		FCDDC6CA212FDF6800E5EF74 /* BatchNormKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCDDC6C9212FDF6800E5EF74 /* BatchNormKernel.metal */; };
-		FCDDC6CC212FDFDB00E5EF74 /* ReluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCDDC6CB212FDFDB00E5EF74 /* ReluKernel.metal */; };
-		FCDDC6CF212FE14700E5EF74 /* PriorBoxKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCDDC6CE212FE14700E5EF74 /* PriorBoxKernel.metal */; };
 		FCDE8A33212A917900F4A8F6 /* ConvTransposeOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCDE8A32212A917900F4A8F6 /* ConvTransposeOp.swift */; };
 		FCE3A1A92153DE5100C37CDE /* ConvAddAddPreluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCE3A1A82153DE5100C37CDE /* ConvAddAddPreluOp.swift */; };
 		FCE3A1AB2153DE8C00C37CDE /* ConvAddAddPreluKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCE3A1AA2153DE8C00C37CDE /* ConvAddAddPreluKernel.swift */; };
 		FCE3A1AD2153E8BA00C37CDE /* ElementwiseAddPreluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCE3A1AC2153E8BA00C37CDE /* ElementwiseAddPreluOp.swift */; };
 		FCE3A1AF2153E8EE00C37CDE /* ElementwiseAddPreluKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCE3A1AE2153E8EE00C37CDE /* ElementwiseAddPreluKernel.swift */; };
-		FCE3A1B12153E90F00C37CDE /* ElementwiseAddPreluKernel.inc.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCE3A1B02153E90F00C37CDE /* ElementwiseAddPreluKernel.inc.metal */; };
-		FCE3A1B32153E91900C37CDE /* ElementwiseAddPreluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCE3A1B22153E91900C37CDE /* ElementwiseAddPreluKernel.metal */; };
 		FCE9D7B7214F869000B520C3 /* Net.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCE9D7B6214F869000B520C3 /* Net.swift */; };
-		FCE9D7B9214FAA4800B520C3 /* NMSFetchResultKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCE9D7B8214FAA4800B520C3 /* NMSFetchResultKernel.metal */; };
-		FCEB684A212F00DB00D2448E /* PreluKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FCEB6849212F00DB00D2448E /* PreluKernel.metal */; };
 		FCEB684C212F093800D2448E /* PreluOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCEB684B212F093800D2448E /* PreluOp.swift */; };
 		FCEBC0F620F1FE120099DBAF /* ConvAddBatchNormReluKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCEBC0F520F1FE120099DBAF /* ConvAddBatchNormReluKernel.swift */; };
 /* End PBXBuildFile section */
@@ -143,34 +105,18 @@
 		456BB7B321F5B356001474E2 /* Framework.pbobjc.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = Framework.pbobjc.h; sourceTree = "<group>"; };
 		4AA1EA852146625E00D0F791 /* BilinearInterpOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = BilinearInterpOp.swift; sourceTree = "<group>"; };
 		4AA1EA87214662BD00D0F791 /* BilinearInterpKernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = BilinearInterpKernel.swift; sourceTree = "<group>"; };
-		4AA1EA892146631C00D0F791 /* BilinearInterp.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BilinearInterp.metal; sourceTree = "<group>"; };
 		4AA1EA8B2146640900D0F791 /* SplitOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = SplitOp.swift; sourceTree = "<group>"; };
 		4AA1EA8D2146647F00D0F791 /* SplitKernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = SplitKernel.swift; sourceTree = "<group>"; };
-		4AA1EA8F214664CD00D0F791 /* Split.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Split.metal; sourceTree = "<group>"; };
 		4AA1EA91214665D700D0F791 /* ShapeOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ShapeOp.swift; sourceTree = "<group>"; };
 		4AA1EA932146661500D0F791 /* ShapeKernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ShapeKernel.swift; sourceTree = "<group>"; };
 		4AA1EA972146666500D0F791 /* FlattenOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = FlattenOp.swift; sourceTree = "<group>"; };
-		4AA1EA9D2148D6F900D0F791 /* ConcatKernel.inc.metal */ = {isa = PBXFileReference; explicitFileType = sourcecode.metal; fileEncoding = 4; path = ConcatKernel.inc.metal; sourceTree = "<group>"; };
-		4AA1EA9F2148DEEE00D0F791 /* ReshapeKernel.inc.metal */ = {isa = PBXFileReference; explicitFileType = sourcecode.metal; fileEncoding = 4; path = ReshapeKernel.inc.metal; sourceTree = "<group>"; };
 		4AA1EAA1214912CC00D0F791 /* FlattenKernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = FlattenKernel.swift; sourceTree = "<group>"; };
-		4AA1EAA3214A295C00D0F791 /* Split.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Split.inc.metal; sourceTree = "<group>"; };
-		4AA1EAA5214B5F6800D0F791 /* Shape.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Shape.metal; sourceTree = "<group>"; };
-		4AA1EAA7214B7AFB00D0F791 /* BilinearInterp.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BilinearInterp.inc.metal; sourceTree = "<group>"; };
-		4AA1EAA9214F53D800D0F791 /* BoxCoder.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BoxCoder.inc.metal; sourceTree = "<group>"; };
-		4AA1EAAB214F55C800D0F791 /* Softmax.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Softmax.inc.metal; sourceTree = "<group>"; };
-		4AA1EAAD214F5FD900D0F791 /* TransposeKernel.inc.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = TransposeKernel.inc.metal; sourceTree = "<group>"; };
-		4AF928762133F1DB005B6C3A /* BoxCoder.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BoxCoder.metal; sourceTree = "<group>"; };
-		4AF9287821341661005B6C3A /* Softmax.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Softmax.metal; sourceTree = "<group>"; };
-		4AF928812135673D005B6C3A /* ConcatKernel.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = ConcatKernel.metal; sourceTree = "<group>"; };
-		4AF9288321357BE3005B6C3A /* Elementwise.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Elementwise.metal; sourceTree = "<group>"; };
 		C28FE02C21BA68C00054EFAC /* Metal.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = Metal.framework; path = System/Library/Frameworks/Metal.framework; sourceTree = SDKROOT; };
 		C28FE02D21BA68C00054EFAC /* MetalPerformanceShaders.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = MetalPerformanceShaders.framework; path = System/Library/Frameworks/MetalPerformanceShaders.framework; sourceTree = SDKROOT; };
 		C28FE02E21BA68C00054EFAC /* MetalKit.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = MetalKit.framework; path = System/Library/Frameworks/MetalKit.framework; sourceTree = SDKROOT; };
 		CDF58151D902A1CBAE56A0C2 /* Pods-paddle-mobile.debug.xcconfig */ = {isa = PBXFileReference; includeInIndex = 1; lastKnownFileType = text.xcconfig; name = "Pods-paddle-mobile.debug.xcconfig"; path = "../Pods/Target Support Files/Pods-paddle-mobile/Pods-paddle-mobile.debug.xcconfig"; sourceTree = "<group>"; };
 		DD2E06330A1E7129C918DB46 /* Pods_paddle_mobile.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; includeInIndex = 0; path = Pods_paddle_mobile.framework; sourceTree = BUILT_PRODUCTS_DIR; };
 		E2A7957C92EDA5C3BEC0FFC2 /* Pods-paddle-mobile.release.xcconfig */ = {isa = PBXFileReference; includeInIndex = 1; lastKnownFileType = text.xcconfig; name = "Pods-paddle-mobile.release.xcconfig"; path = "../Pods/Target Support Files/Pods-paddle-mobile/Pods-paddle-mobile.release.xcconfig"; sourceTree = "<group>"; };
-		FC0226552138F33800F395E2 /* TransposeKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = TransposeKernel.metal; sourceTree = "<group>"; };
-		FC0226572138F38D00F395E2 /* PoolKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = PoolKernel.metal; sourceTree = "<group>"; };
 		FC039B6A20E11C3C0081E9F8 /* paddle_mobile.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; includeInIndex = 0; path = paddle_mobile.framework; sourceTree = BUILT_PRODUCTS_DIR; };
 		FC039B6D20E11C3C0081E9F8 /* paddle_mobile.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = paddle_mobile.h; sourceTree = "<group>"; };
 		FC039B6E20E11C3C0081E9F8 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; sourceTree = "<group>"; };
@@ -198,41 +144,26 @@
 		FC0E2DBB20EE45FE009C1FAC /* ConvKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvKernel.swift; sourceTree = "<group>"; };
 		FC0E2DBD20EE460D009C1FAC /* BatchNormKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = BatchNormKernel.swift; sourceTree = "<group>"; };
 		FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ElementwiseAddKernel.swift; sourceTree = "<group>"; };
-		FC1B16B220EC9A4F00678B91 /* Kernels.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = Kernels.metal; sourceTree = "<group>"; };
 		FC1CF3F621D4B4C400F7392E /* Runner.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Runner.swift; sourceTree = "<group>"; };
 		FC2BFCC121DF2F9100C262B2 /* GlobalConfig.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = GlobalConfig.swift; sourceTree = "<group>"; };
 		FC2BFD4521DF685F00C262B2 /* Scale.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Scale.swift; sourceTree = "<group>"; };
 		FC2BFD4921DF81DE00C262B2 /* Kernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = Kernel.swift; sourceTree = "<group>"; };
 		FC2BFD4D21DF820A00C262B2 /* ConvAddBatchNormReluOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ConvAddBatchNormReluOp.swift; sourceTree = "<group>"; };
-		FC2BFD5021DF8E0400C262B2 /* Scale.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = Scale.metal; sourceTree = "<group>"; };
 		FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PaddleMobileUnitTest.swift; sourceTree = "<group>"; };
-		FC4CB74820F0B954007C0C6D /* ConvKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvKernel.metal; sourceTree = "<group>"; };
 		FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ProgramOptimize.swift; sourceTree = "<group>"; };
 		FC4FD97D2140F2C30073E130 /* libstdc++.tbd */ = {isa = PBXFileReference; lastKnownFileType = "sourcecode.text-based-dylib-definition"; name = "libstdc++.tbd"; path = "usr/lib/libstdc++.tbd"; sourceTree = SDKROOT; };
 		FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Texture2DTo2DArrayKernel.swift; sourceTree = "<group>"; };
 		FC60DB8820E9AAA500FF203F /* MetalExtension.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MetalExtension.swift; sourceTree = "<group>"; };
 		FC803BBE214CB65A0094B8E5 /* ConvAddPreluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvAddPreluOp.swift; sourceTree = "<group>"; };
 		FC803BC0214CB77A0094B8E5 /* ConvAddPreluKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvAddPreluKernel.swift; sourceTree = "<group>"; };
-		FC803BC2214CB79C0094B8E5 /* ConvAddPreluKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvAddPreluKernel.metal; sourceTree = "<group>"; };
-		FC803BC4214CB8F00094B8E5 /* ConvAddPrelu.inc.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvAddPrelu.inc.metal; sourceTree = "<group>"; };
-		FC803BC6214CBA820094B8E5 /* Macro.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = Macro.metal; sourceTree = "<group>"; };
-		FC803BC8214CFC8D0094B8E5 /* FetchKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = FetchKernel.metal; sourceTree = "<group>"; };
 		FC82735820E3C04200BE430A /* OpCreator.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpCreator.swift; sourceTree = "<group>"; };
 		FC9797C821D6101D00F2FD90 /* ResizeBilinearOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ResizeBilinearOp.swift; sourceTree = "<group>"; };
 		FC9797CA21D6102D00F2FD90 /* ResizeBilinearKernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ResizeBilinearKernel.swift; sourceTree = "<group>"; };
-		FC9C2A0C21D3D185005856C6 /* FetchKernel.inc.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = FetchKernel.inc.metal; sourceTree = "<group>"; };
 		FC9D037820E229E4000F735A /* OpParam.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpParam.swift; sourceTree = "<group>"; };
 		FC9D037F20E22FBB000F735A /* FeedOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FeedOp.swift; sourceTree = "<group>"; };
 		FC9D038120E2312E000F735A /* FetchOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FetchOp.swift; sourceTree = "<group>"; };
 		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>"; };
-		FCA67B1621364EF000BD58AA /* ConvTransposeKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvTransposeKernel.metal; sourceTree = "<group>"; };
-		FCA67CD42138272900BD58AA /* ConvAddMetal.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvAddMetal.metal; sourceTree = "<group>"; };
-		FCA67CD6213827AC00BD58AA /* ConvAddBNReluKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvAddBNReluKernel.metal; sourceTree = "<group>"; };
-		FCA67CD82138287B00BD58AA /* ConvBNReluKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvBNReluKernel.metal; sourceTree = "<group>"; };
 		FCB40E5821E0DCAB0075EC91 /* FetchKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FetchKernel.swift; sourceTree = "<group>"; };
-		FCB91DC121FEEE990051C6B2 /* BufferToTexture.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = BufferToTexture.metal; sourceTree = "<group>"; };
 		FCBCCC562122F41300D94F7E /* DwConvBNReluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = DwConvBNReluOp.swift; sourceTree = "<group>"; };
 		FCBCCC582122F42700D94F7E /* ConvBNReluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvBNReluOp.swift; sourceTree = "<group>"; };
 		FCBCCC5A2122F66F00D94F7E /* ConvBNReluKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvBNReluKernel.swift; sourceTree = "<group>"; };
@@ -247,7 +178,6 @@
 		FCBCCC6C2123073A00D94F7E /* BoxcoderKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = BoxcoderKernel.swift; sourceTree = "<group>"; };
 		FCBCCC6E2123097100D94F7E /* MulticlassNMSOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MulticlassNMSOp.swift; sourceTree = "<group>"; };
 		FCBCCC70212309A700D94F7E /* MulticlassNMSKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MulticlassNMSKernel.swift; sourceTree = "<group>"; };
-		FCCED5E021D71FC000BE8D5F /* PoolKernel.inc.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = PoolKernel.inc.metal; sourceTree = "<group>"; };
 		FCD04E6520F314C50007374F /* PoolOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PoolOp.swift; sourceTree = "<group>"; };
 		FCD04E6720F315020007374F /* PoolKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PoolKernel.swift; sourceTree = "<group>"; };
 		FCD04E6920F319EC0007374F /* SoftmaxOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = SoftmaxOp.swift; sourceTree = "<group>"; };
@@ -259,19 +189,12 @@
 		FCDC0FEA21099A1D00DC9EFB /* Tools.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Tools.swift; sourceTree = "<group>"; };
 		FCDDC6C5212F9FB800E5EF74 /* PreluKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PreluKernel.swift; sourceTree = "<group>"; };
 		FCDDC6C7212FA3CA00E5EF74 /* ConvTransposeKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvTransposeKernel.swift; sourceTree = "<group>"; };
-		FCDDC6C9212FDF6800E5EF74 /* BatchNormKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = BatchNormKernel.metal; sourceTree = "<group>"; };
-		FCDDC6CB212FDFDB00E5EF74 /* ReluKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ReluKernel.metal; sourceTree = "<group>"; };
-		FCDDC6CE212FE14700E5EF74 /* PriorBoxKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = PriorBoxKernel.metal; sourceTree = "<group>"; };
 		FCDE8A32212A917900F4A8F6 /* ConvTransposeOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvTransposeOp.swift; sourceTree = "<group>"; };
 		FCE3A1A82153DE5100C37CDE /* ConvAddAddPreluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvAddAddPreluOp.swift; sourceTree = "<group>"; };
 		FCE3A1AA2153DE8C00C37CDE /* ConvAddAddPreluKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvAddAddPreluKernel.swift; sourceTree = "<group>"; };
 		FCE3A1AC2153E8BA00C37CDE /* ElementwiseAddPreluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ElementwiseAddPreluOp.swift; sourceTree = "<group>"; };
 		FCE3A1AE2153E8EE00C37CDE /* ElementwiseAddPreluKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ElementwiseAddPreluKernel.swift; sourceTree = "<group>"; };
-		FCE3A1B02153E90F00C37CDE /* ElementwiseAddPreluKernel.inc.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ElementwiseAddPreluKernel.inc.metal; sourceTree = "<group>"; };
-		FCE3A1B22153E91900C37CDE /* ElementwiseAddPreluKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ElementwiseAddPreluKernel.metal; sourceTree = "<group>"; };
 		FCE9D7B6214F869000B520C3 /* Net.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Net.swift; sourceTree = "<group>"; };
-		FCE9D7B8214FAA4800B520C3 /* NMSFetchResultKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = NMSFetchResultKernel.metal; sourceTree = "<group>"; };
-		FCEB6849212F00DB00D2448E /* PreluKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = PreluKernel.metal; sourceTree = "<group>"; };
 		FCEB684B212F093800D2448E /* PreluOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PreluOp.swift; sourceTree = "<group>"; };
 		FCEBC0F520F1FE120099DBAF /* ConvAddBatchNormReluKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ConvAddBatchNormReluKernel.swift; sourceTree = "<group>"; };
 /* End PBXFileReference section */
@@ -426,7 +349,6 @@
 			isa = PBXGroup;
 			children = (
 				FCDDC6CD212FE02100E5EF74 /* Base */,
-				FCEB6837212F00B100D2448E /* metal */,
 				FC9797CA21D6102D00F2FD90 /* ResizeBilinearKernel.swift */,
 				FCDDC6C7212FA3CA00E5EF74 /* ConvTransposeKernel.swift */,
 				FC0E2DBB20EE45FE009C1FAC /* ConvKernel.swift */,
@@ -498,52 +420,6 @@
 			path = Base;
 			sourceTree = "<group>";
 		};
-		FCEB6837212F00B100D2448E /* metal */ = {
-			isa = PBXGroup;
-			children = (
-				FCB91DC121FEEE990051C6B2 /* BufferToTexture.metal */,
-				4AF928812135673D005B6C3A /* ConcatKernel.metal */,
-				4AA1EA9D2148D6F900D0F791 /* ConcatKernel.inc.metal */,
-				4AF9288321357BE3005B6C3A /* Elementwise.metal */,
-				FC1B16B220EC9A4F00678B91 /* Kernels.metal */,
-				FC4CB74820F0B954007C0C6D /* ConvKernel.metal */,
-				4AF928762133F1DB005B6C3A /* BoxCoder.metal */,
-				4AA1EAA9214F53D800D0F791 /* BoxCoder.inc.metal */,
-				4AA1EAA5214B5F6800D0F791 /* Shape.metal */,
-				4AA1EA8F214664CD00D0F791 /* Split.metal */,
-				4AA1EAA3214A295C00D0F791 /* Split.inc.metal */,
-				4AA1EA892146631C00D0F791 /* BilinearInterp.metal */,
-				4AA1EAA7214B7AFB00D0F791 /* BilinearInterp.inc.metal */,
-				4AF9287821341661005B6C3A /* Softmax.metal */,
-				4AA1EAAB214F55C800D0F791 /* Softmax.inc.metal */,
-				FCEB6849212F00DB00D2448E /* PreluKernel.metal */,
-				FCDDC6C9212FDF6800E5EF74 /* BatchNormKernel.metal */,
-				FCDDC6CB212FDFDB00E5EF74 /* ReluKernel.metal */,
-				FCDDC6CE212FE14700E5EF74 /* PriorBoxKernel.metal */,
-				FCA3A1622132A4AC00084FE5 /* ReshapeKernel.metal */,
-				4AA1EA9F2148DEEE00D0F791 /* ReshapeKernel.inc.metal */,
-				FCA3A1642132A5EB00084FE5 /* Common.metal */,
-				FCA67B1621364EF000BD58AA /* ConvTransposeKernel.metal */,
-				FCA67CD42138272900BD58AA /* ConvAddMetal.metal */,
-				FCA67CD6213827AC00BD58AA /* ConvAddBNReluKernel.metal */,
-				FCA67CD82138287B00BD58AA /* ConvBNReluKernel.metal */,
-				FC0226552138F33800F395E2 /* TransposeKernel.metal */,
-				4AA1EAAD214F5FD900D0F791 /* TransposeKernel.inc.metal */,
-				FC0226572138F38D00F395E2 /* PoolKernel.metal */,
-				FCCED5E021D71FC000BE8D5F /* PoolKernel.inc.metal */,
-				FC803BC2214CB79C0094B8E5 /* ConvAddPreluKernel.metal */,
-				FC803BC4214CB8F00094B8E5 /* ConvAddPrelu.inc.metal */,
-				FC803BC6214CBA820094B8E5 /* Macro.metal */,
-				FC803BC8214CFC8D0094B8E5 /* FetchKernel.metal */,
-				FC9C2A0C21D3D185005856C6 /* FetchKernel.inc.metal */,
-				FCE9D7B8214FAA4800B520C3 /* NMSFetchResultKernel.metal */,
-				FCE3A1B02153E90F00C37CDE /* ElementwiseAddPreluKernel.inc.metal */,
-				FCE3A1B22153E91900C37CDE /* ElementwiseAddPreluKernel.metal */,
-				FC2BFD5021DF8E0400C262B2 /* Scale.metal */,
-			);
-			path = metal;
-			sourceTree = "<group>";
-		};
 /* End PBXGroup section */
 
 /* Begin PBXHeadersBuildPhase section */
@@ -647,67 +523,41 @@
 			buildActionMask = 2147483647;
 			files = (
 				FC9D038020E22FBB000F735A /* FeedOp.swift in Sources */,
-				FC9C2A0D21D3D185005856C6 /* FetchKernel.inc.metal in Sources */,
-				4AA1EAAA214F53D800D0F791 /* BoxCoder.inc.metal in Sources */,
 				FC039B9F20E11CB20081E9F8 /* Tensor.swift in Sources */,
-				FC803BC9214CFC8D0094B8E5 /* FetchKernel.metal in Sources */,
-				FCA67CD7213827AC00BD58AA /* ConvAddBNReluKernel.metal in Sources */,
-				4AF9287921341661005B6C3A /* Softmax.metal in Sources */,
 				4AA1EA942146661500D0F791 /* ShapeKernel.swift in Sources */,
 				FC0E2DBC20EE45FE009C1FAC /* ConvKernel.swift in Sources */,
 				FC039BAA20E11CBC0081E9F8 /* ElementwiseAddOp.swift in Sources */,
 				FCDE8A33212A917900F4A8F6 /* ConvTransposeOp.swift in Sources */,
 				FCBCCC6B2123071700D94F7E /* BoxcoderOp.swift in Sources */,
-				4AA1EAAE214F5FD900D0F791 /* TransposeKernel.inc.metal in Sources */,
-				4AA1EAA4214A295C00D0F791 /* Split.inc.metal in Sources */,
-				FC803BC7214CBA820094B8E5 /* Macro.metal in Sources */,
 				FC039B9B20E11CA00081E9F8 /* Executor.swift in Sources */,
-				4AF9288421357BE3005B6C3A /* Elementwise.metal in Sources */,
 				FCD04E7020F31B720007374F /* ReshapeKernel.swift in Sources */,
-				FCE3A1B12153E90F00C37CDE /* ElementwiseAddPreluKernel.inc.metal in Sources */,
 				FCD04E7220F343420007374F /* ConvAddOp.swift in Sources */,
 				FC039BBB20E11CC20081E9F8 /* PMProgramDesc.swift in Sources */,
 				FCE3A1AB2153DE8C00C37CDE /* ConvAddAddPreluKernel.swift in Sources */,
 				FC9D037920E229E4000F735A /* OpParam.swift in Sources */,
 				FC3602CC2108819F00FACB58 /* PaddleMobileUnitTest.swift in Sources */,
-				FCDDC6CC212FDFDB00E5EF74 /* ReluKernel.metal in Sources */,
-				FC0226562138F33800F395E2 /* TransposeKernel.metal in Sources */,
 				FCDDC6C6212F9FB800E5EF74 /* PreluKernel.swift in Sources */,
 				FC9797CB21D6102D00F2FD90 /* ResizeBilinearKernel.swift in Sources */,
-				FCA67CD52138272900BD58AA /* ConvAddMetal.metal in Sources */,
 				FCBCCC5B2122F66F00D94F7E /* ConvBNReluKernel.swift in Sources */,
 				4AA1EA8C2146640900D0F791 /* SplitOp.swift in Sources */,
-				FCB91DC221FEEE990051C6B2 /* BufferToTexture.metal in Sources */,
-				4AA1EAAC214F55C800D0F791 /* Softmax.inc.metal in Sources */,
 				FC0E2DC020EE461F009C1FAC /* ElementwiseAddKernel.swift in Sources */,
-				4AF928772133F1DB005B6C3A /* BoxCoder.metal in Sources */,
 				FC803BBF214CB65A0094B8E5 /* ConvAddPreluOp.swift in Sources */,
 				FCEB684C212F093800D2448E /* PreluOp.swift in Sources */,
-				4AA1EAA8214B7AFB00D0F791 /* BilinearInterp.inc.metal in Sources */,
-				FCA67CD92138287B00BD58AA /* ConvBNReluKernel.metal in Sources */,
 				FC60DB8920E9AAA500FF203F /* MetalExtension.swift in Sources */,
 				FCEBC0F620F1FE120099DBAF /* ConvAddBatchNormReluKernel.swift in Sources */,
-				4AA1EA8A2146631C00D0F791 /* BilinearInterp.metal in Sources */,
-				FCDDC6CA212FDF6800E5EF74 /* BatchNormKernel.metal in Sources */,
-				FC1B16B320EC9A4F00678B91 /* Kernels.metal in Sources */,
 				FC039BBA20E11CC20081E9F8 /* TensorDesc.swift in Sources */,
 				FC039BA020E11CB20081E9F8 /* Dim.swift in Sources */,
 				FC039B9920E11C9A0081E9F8 /* Types.swift in Sources */,
-				FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */,
-				FCA3A1632132A4AC00084FE5 /* ReshapeKernel.metal in Sources */,
 				FCBCCC592122F42700D94F7E /* ConvBNReluOp.swift in Sources */,
 				FC039BA920E11CBC0081E9F8 /* ConvOp.swift in Sources */,
-				FCCED5E121D71FC000BE8D5F /* PoolKernel.inc.metal in Sources */,
 				FC2BFD4A21DF81DE00C262B2 /* Kernel.swift in Sources */,
 				FC9D038420E23B01000F735A /* Texture.swift in Sources */,
-				FCE3A1B32153E91900C37CDE /* ElementwiseAddPreluKernel.metal in Sources */,
 				FC2BFD4E21DF820B00C262B2 /* ConvAddBatchNormReluOp.swift in Sources */,
 				4AA1EAA2214912CD00D0F791 /* FlattenKernel.swift in Sources */,
 				4AA1EA982146666500D0F791 /* FlattenOp.swift in Sources */,
 				FC2BFCC221DF2F9100C262B2 /* GlobalConfig.swift in Sources */,
 				456BB7B421F5B356001474E2 /* Framework.pbobjc.m in Sources */,
 				FCBCCC652122FCD700D94F7E /* TransposeOp.swift in Sources */,
-				4AA1EAA6214B5F6800D0F791 /* Shape.metal in Sources */,
 				FCD04E6E20F31B4B0007374F /* ReshapeOp.swift in Sources */,
 				FC039B9820E11C9A0081E9F8 /* Errors.swift in Sources */,
 				FC039BBF20E11CC20081E9F8 /* Attribute.swift in Sources */,
@@ -716,12 +566,9 @@
 				FC1CF3F721D4B4C400F7392E /* Runner.swift in Sources */,
 				FC039BB920E11CC20081E9F8 /* Scope.swift in Sources */,
 				FCD04E6620F314C50007374F /* PoolOp.swift in Sources */,
-				FCE9D7B9214FAA4800B520C3 /* NMSFetchResultKernel.metal in Sources */,
 				FC039BAC20E11CBC0081E9F8 /* BatchNormOp.swift in Sources */,
 				FCBCCC6F2123097100D94F7E /* MulticlassNMSOp.swift in Sources */,
 				FC039BBC20E11CC20081E9F8 /* PMVarDesc.swift in Sources */,
-				FC803BC5214CB8F00094B8E5 /* ConvAddPrelu.inc.metal in Sources */,
-				4AF928822135673D005B6C3A /* ConcatKernel.metal in Sources */,
 				FCBCCC632122FCC000D94F7E /* TransposeKernel.swift in Sources */,
 				FCBCCC71212309A700D94F7E /* MulticlassNMSKernel.swift in Sources */,
 				FCDC0FEB21099A1D00DC9EFB /* Tools.swift in Sources */,
@@ -732,7 +579,6 @@
 				FCBCCC69212306D300D94F7E /* ConcatKernel.swift in Sources */,
 				FCDDC6C8212FA3CA00E5EF74 /* ConvTransposeKernel.swift in Sources */,
 				FC82735920E3C04200BE430A /* OpCreator.swift in Sources */,
-				FCA3A1652132A5EB00084FE5 /* Common.metal in Sources */,
 				4AA1EA92214665D700D0F791 /* ShapeOp.swift in Sources */,
 				FC803BC1214CB77A0094B8E5 /* ConvAddPreluKernel.swift in Sources */,
 				FCBCCC5D2122F8A100D94F7E /* DepthwiseConvOp.swift in Sources */,
@@ -744,24 +590,16 @@
 				FCBCCC612122FBDF00D94F7E /* PriorBoxKernel.swift in Sources */,
 				FCBCCC5F2122FB3B00D94F7E /* PriorBoxOp.swift in Sources */,
 				FC9D038220E2312E000F735A /* FetchOp.swift in Sources */,
-				FCA67B1721364EF000BD58AA /* ConvTransposeKernel.metal in Sources */,
 				FC039BBD20E11CC20081E9F8 /* Program.swift in Sources */,
-				FC2BFD5121DF8E0400C262B2 /* Scale.metal in Sources */,
 				FC039BA220E11CB70081E9F8 /* Loader.swift in Sources */,
 				FCBCCC67212306B000D94F7E /* ConcatOp.swift in Sources */,
 				FCD04E6C20F31A280007374F /* SoftmaxKernel.swift in Sources */,
-				FCEB684A212F00DB00D2448E /* PreluKernel.metal in Sources */,
-				4AA1EAA02148DEEE00D0F791 /* ReshapeKernel.inc.metal in Sources */,
-				FCDDC6CF212FE14700E5EF74 /* PriorBoxKernel.metal in Sources */,
 				FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */,
 				FCE3A1A92153DE5100C37CDE /* ConvAddAddPreluOp.swift in Sources */,
 				FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */,
 				FCE3A1AD2153E8BA00C37CDE /* ElementwiseAddPreluOp.swift in Sources */,
 				FC039BC020E11CC20081E9F8 /* PMBlockDesc.swift in Sources */,
-				FC803BC3214CB79C0094B8E5 /* ConvAddPreluKernel.metal in Sources */,
-				4AA1EA90214664CD00D0F791 /* Split.metal in Sources */,
 				FCD04E6820F315020007374F /* PoolKernel.swift in Sources */,
-				FC0226582138F38D00F395E2 /* PoolKernel.metal in Sources */,
 				FC039BAD20E11CBC0081E9F8 /* ReluOp.swift in Sources */,
 				FCBCCC572122F41300D94F7E /* DwConvBNReluOp.swift in Sources */,
 				FC039BBE20E11CC20081E9F8 /* PMOpDesc.swift in Sources */,

metal/paddle-mobile/paddle-mobile/API/Net.swift

@@ -57,11 +57,11 @@ import Foundation
   @objc public var inputDim: Dim = Dim.init(inDim: [])
   
   
-  @objc public init(device: MTLDevice, paramPointer: UnsafeMutableRawPointer, paramSize:Int, modePointer: UnsafeMutableRawPointer, modelSize: Int) {
-    self.paramPointer = paramPointer
-    self.paramSize = paramSize
-    self.modelPointer = modePointer
-    self.modelSize = modelSize
+  @objc public init(device: MTLDevice, inParamPointer: UnsafeMutableRawPointer, inParamSize:Int, inModelPointer: UnsafeMutableRawPointer, inModelSize: Int) {
+    self.paramPointer = inParamPointer
+    self.paramSize = inParamSize
+    self.modelPointer = inModelPointer
+    self.modelSize = inModelSize
     self.device = device
     super.init()
   }
@@ -80,7 +80,7 @@ import Foundation
       guard let inResPointer = gpuRes.resultPointer else {
         fatalError()
       }
-      return ResultHolder.init(inResult: inResPointer, inCapacity: paddleMobileRes.capacity, inDim: gpuRes.dim)
+      return ResultHolder.init(inResult: inResPointer, inCapacity: gpuRes.capacity, inDim: gpuRes.dim)
     }
   }
   

src/framework/loader.cpp

@@ -241,6 +241,7 @@ const Program<Device, T> Loader<Device, T>::LoadProgram(
   FusionAndPrintInfos(optimize, can_add_split, &program, originProgramDesc);
 
   paddle_mobile__framework__proto__program_desc__free_unpacked(c_program, NULL);
+  free(buf);
   return program;
 }
 

src/io/ios_io/PaddleMobileCPU.mm

@@ -96,6 +96,7 @@ static std::mutex shared_mutex;
 - (void)dealloc {
   if (pam_) {
     delete pam_;
+    pam_ = nullptr;
   }
 }