update paddle mobile demo

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

@@ -18,17 +18,17 @@
 		FC27991321343A3A000B6BAD /* CPUCompute.mm in Sources */ = {isa = PBXBuildFile; fileRef = FC27991221343A3A000B6BAD /* CPUCompute.mm */; };
 		FC3C800F2133F46600D1295E /* MobileNetSSD.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3C800E2133F46600D1295E /* MobileNetSSD.swift */; };
 		FC3C80112133F4AB00D1295E /* MobileNet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3C80102133F4AB00D1295E /* MobileNet.swift */; };
-		FC8CFEDF213521C10094D569 /* genet_model in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEDD213521C10094D569 /* genet_model */; };
-		FC8CFEE0213521C10094D569 /* genet_params in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEDE213521C10094D569 /* genet_params */; };
 		FC8CFEE2213524EA0094D569 /* Genet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC8CFEE1213524EA0094D569 /* Genet.swift */; };
+		FC8CFEE62135452C0094D569 /* genet_params in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEE42135452B0094D569 /* genet_params */; };
+		FC8CFEE72135452C0094D569 /* genet_model in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEE52135452B0094D569 /* genet_model */; };
+		FC8CFEF8213551D10094D569 /* params in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEF6213551D00094D569 /* params */; };
+		FC8CFEF9213551D10094D569 /* model in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEF7213551D00094D569 /* model */; };
 		FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */ = {isa = PBXBuildFile; fileRef = FC918190211DBC3500B6F354 /* paddle-mobile.png */; };
 		FC918193211DC70500B6F354 /* iphone.JPG in Resources */ = {isa = PBXBuildFile; fileRef = FC918192211DC70500B6F354 /* iphone.JPG */; };
 		FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FCA3A16021313E1F00084FE5 /* hand.jpg */; };
 		FCBCCC522122EEDC00D94F7E /* ssd_hand_params in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC502122EEDC00D94F7E /* ssd_hand_params */; };
 		FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC512122EEDC00D94F7E /* ssd_hand_model */; };
 		FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC542122EF5400D94F7E /* MetalHelper.swift */; };
-		FCD04E6320F3146B0007374F /* params in Resources */ = {isa = PBXBuildFile; fileRef = FCD04E6120F3146A0007374F /* params */; };
-		FCD04E6420F3146B0007374F /* model in Resources */ = {isa = PBXBuildFile; fileRef = FCD04E6220F3146A0007374F /* model */; };
 		FCDFD41B211D91C7005AB38B /* synset.txt in Resources */ = {isa = PBXBuildFile; fileRef = FCDFD41A211D91C7005AB38B /* synset.txt */; };
 		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, ); }; };
@@ -67,17 +67,17 @@
 		FC27991421343A46000B6BAD /* CPUCompute.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = CPUCompute.h; sourceTree = "<group>"; };
 		FC3C800E2133F46600D1295E /* MobileNetSSD.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNetSSD.swift; sourceTree = "<group>"; };
 		FC3C80102133F4AB00D1295E /* MobileNet.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNet.swift; sourceTree = "<group>"; };
-		FC8CFEDD213521C10094D569 /* genet_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = genet_model; sourceTree = "<group>"; };
-		FC8CFEDE213521C10094D569 /* genet_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = genet_params; sourceTree = "<group>"; };
 		FC8CFEE1213524EA0094D569 /* Genet.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Genet.swift; sourceTree = "<group>"; };
+		FC8CFEE42135452B0094D569 /* genet_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = genet_params; sourceTree = "<group>"; };
+		FC8CFEE52135452B0094D569 /* genet_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = genet_model; sourceTree = "<group>"; };
+		FC8CFEF6213551D00094D569 /* params */ = {isa = PBXFileReference; lastKnownFileType = file; path = params; sourceTree = "<group>"; };
+		FC8CFEF7213551D00094D569 /* model */ = {isa = PBXFileReference; lastKnownFileType = file; path = model; sourceTree = "<group>"; };
 		FC918190211DBC3500B6F354 /* paddle-mobile.png */ = {isa = PBXFileReference; lastKnownFileType = image.png; path = "paddle-mobile.png"; sourceTree = "<group>"; };
 		FC918192211DC70500B6F354 /* iphone.JPG */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = iphone.JPG; sourceTree = "<group>"; };
 		FCA3A16021313E1F00084FE5 /* hand.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = hand.jpg; sourceTree = "<group>"; };
 		FCBCCC502122EEDC00D94F7E /* ssd_hand_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_params; sourceTree = "<group>"; };
 		FCBCCC512122EEDC00D94F7E /* ssd_hand_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_model; sourceTree = "<group>"; };
 		FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = "<group>"; };
-		FCD04E6120F3146A0007374F /* params */ = {isa = PBXFileReference; lastKnownFileType = file; path = params; sourceTree = "<group>"; };
-		FCD04E6220F3146A0007374F /* model */ = {isa = PBXFileReference; lastKnownFileType = file; path = model; sourceTree = "<group>"; };
 		FCDFD41A211D91C7005AB38B /* synset.txt */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = synset.txt; sourceTree = "<group>"; };
 		FCEBEC2B20E1391F00C0B14D /* paddle_mobile.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; path = paddle_mobile.framework; sourceTree = BUILT_PRODUCTS_DIR; };
 		FCEEE7D3210627A000444BEC /* banana.jpeg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = banana.jpeg; sourceTree = "<group>"; };
@@ -165,9 +165,9 @@
 		FC0E2C2020EDC03B009C1FAC /* models */ = {
 			isa = PBXGroup;
 			children = (
-				FC8CFED821351F5D0094D569 /* genet */,
+				FC8CFEF5213551D00094D569 /* mobilenet */,
+				FC8CFEE32135452B0094D569 /* genet */,
 				FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */,
-				FCD04E6020F3146A0007374F /* mobilenet */,
 			);
 			name = models;
 			path = ../../models;
@@ -188,31 +188,31 @@
 			path = Net;
 			sourceTree = "<group>";
 		};
-		FC8CFED821351F5D0094D569 /* genet */ = {
+		FC8CFEE32135452B0094D569 /* genet */ = {
 			isa = PBXGroup;
 			children = (
-				FC8CFEDD213521C10094D569 /* genet_model */,
-				FC8CFEDE213521C10094D569 /* genet_params */,
+				FC8CFEE42135452B0094D569 /* genet_params */,
+				FC8CFEE52135452B0094D569 /* genet_model */,
 			);
 			path = genet;
 			sourceTree = "<group>";
 		};
-		FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */ = {
+		FC8CFEF5213551D00094D569 /* mobilenet */ = {
 			isa = PBXGroup;
 			children = (
-				FCBCCC502122EEDC00D94F7E /* ssd_hand_params */,
-				FCBCCC512122EEDC00D94F7E /* ssd_hand_model */,
+				FC8CFEF6213551D00094D569 /* params */,
+				FC8CFEF7213551D00094D569 /* model */,
 			);
-			path = mobilenet_ssd_hand;
+			path = mobilenet;
 			sourceTree = "<group>";
 		};
-		FCD04E6020F3146A0007374F /* mobilenet */ = {
+		FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */ = {
 			isa = PBXGroup;
 			children = (
-				FCD04E6120F3146A0007374F /* params */,
-				FCD04E6220F3146A0007374F /* model */,
+				FCBCCC502122EEDC00D94F7E /* ssd_hand_params */,
+				FCBCCC512122EEDC00D94F7E /* ssd_hand_model */,
 			);
-			path = mobilenet;
+			path = mobilenet_ssd_hand;
 			sourceTree = "<group>";
 		};
 /* End PBXGroup section */
@@ -277,19 +277,19 @@
 			isa = PBXResourcesBuildPhase;
 			buildActionMask = 2147483647;
 			files = (
-				FCD04E6320F3146B0007374F /* params in Resources */,
+				FC8CFEF8213551D10094D569 /* params in Resources */,
 				FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */,
-				FC8CFEE0213521C10094D569 /* genet_params in Resources */,
+				FC8CFEF9213551D10094D569 /* model in Resources */,
 				FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */,
+				FC8CFEE72135452C0094D569 /* genet_model in Resources */,
 				FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */,
 				FCBCCC522122EEDC00D94F7E /* ssd_hand_params in Resources */,
 				FCEEE7D4210627A000444BEC /* banana.jpeg in Resources */,
 				FC918193211DC70500B6F354 /* iphone.JPG in Resources */,
 				FCDFD41B211D91C7005AB38B /* synset.txt in Resources */,
-				FCD04E6420F3146B0007374F /* model in Resources */,
-				FC8CFEDF213521C10094D569 /* genet_model in Resources */,
 				FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */,
 				FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */,
+				FC8CFEE62135452C0094D569 /* genet_params in Resources */,
 				FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */,
 			);
 			runOnlyForDeploymentPostprocessing = 0;

metal/paddle-mobile-demo/paddle-mobile-demo/Net/CPUCompute.mm

@@ -317,18 +317,3 @@ void MultiClassNMSCompute(NMSParam *param) {
 @end
 
 
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-

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

@@ -35,7 +35,7 @@ class Genet: Net {
   }
   
   var preprocessKernel: CusomKernel
-  let dim = [1, 128, 128, 3]
+  let dim = (n: 1, h: 128, w: 128, c: 3)
   let modelPath: String
   let paramPath: String
   let modelDir: String

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

-//
-//  MetalHelper.swift
-//  paddle-mobile-demo
-//
-//  Created by liuRuiLong on 2018/7/25.
-//  Copyright © 2018年 orange. All rights reserved.
-//
+/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
+ 
+ Licensed under the Apache License, Version 2.0 (the "License");
+ you may not use this file except in compliance with the License.
+ You may obtain a copy of the License at
+ 
+ http://www.apache.org/licenses/LICENSE-2.0
+ 
+ Unless required by applicable law or agreed to in writing, software
+ distributed under the License is distributed on an "AS IS" BASIS,
+ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ See the License for the specific language governing permissions and
+ limitations under the License. */
 
 import Metal
 import MetalKit
 import Foundation
 import paddle_mobile
-import MetalPerformanceShaders
 
 class MetalHelper {
-    let device: MTLDevice
-    let queue: MTLCommandQueue
-    let textureLoader: MTKTextureLoader
-    static let shared: MetalHelper = MetalHelper.init()
-    private init(){
-        device = MTLCreateSystemDefaultDevice()!
-        queue = device.makeCommandQueue()!
-        textureLoader = MTKTextureLoader.init(device: device)
+  let device: MTLDevice
+  let queue: MTLCommandQueue
+  let textureLoader: MTKTextureLoader
+  static let shared: MetalHelper = MetalHelper.init()
+  private init(){
+    device = MTLCreateSystemDefaultDevice()!
+    queue = device.makeCommandQueue()!
+    textureLoader = MTKTextureLoader.init(device: device)
+  }
+  
+  static func scaleTexture(queue: MTLCommandQueue, input: MTLTexture, size:(width: Int, height: Int), complete: @escaping (MTLTexture) -> Void) {
+
+    guard let buffer = queue.makeCommandBuffer() else {
+      fatalError()
+    }
+    
+    let scaleKernel = ScaleKernel.init(device: MetalHelper.shared.device, shape: CusomKernel.Shape.init(inWidth: size.width, inHeight: size.height, inChannel: 3))
+    
+    do {
+      try scaleKernel.compute(inputTexuture: input, commandBuffer: buffer)
+    } catch let error {
+      print(error)
+      fatalError()
     }
     
-    static func scaleTexture(queue: MTLCommandQueue, input: MTLTexture, size:(width: Int, height: Int), complete: @escaping (MTLTexture) -> Void) {
-        let tmpTextureDes = MTLTextureDescriptor.init()
-        tmpTextureDes.width = size.width
-        tmpTextureDes.height = size.height
-        tmpTextureDes.depth = 1
-        tmpTextureDes.usage = [.shaderRead, .shaderWrite]
-        tmpTextureDes.pixelFormat = .rgba32Float
-        tmpTextureDes.textureType = .type2D
-        tmpTextureDes.storageMode = .shared
-        tmpTextureDes.cpuCacheMode = .defaultCache
-        let dest = MetalHelper.shared.device.makeTexture(descriptor: tmpTextureDes)
-        
-        let scale = MPSImageLanczosScale.init(device: MetalHelper.shared.device)
-        
-        let buffer = queue.makeCommandBuffer()
-        scale.encode(commandBuffer: buffer!, sourceTexture: input, destinationTexture: dest!)
-        buffer?.addCompletedHandler({ (buffer) in
-            complete(dest!)
-        })
-        buffer?.commit()
+    buffer.addCompletedHandler { (buffer) in
+      complete(scaleKernel.outputTexture)
     }
+    buffer.commit()
+  }
 }
 

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

@@ -58,7 +58,7 @@ class MobileNet: Net{
   }
   
   var preprocessKernel: CusomKernel
-  let dim = [1, 224, 224, 3]
+  let dim = (n: 1, h: 224, w: 224, c: 3)
   let modelPath: String
   let paramPath: String
   let modelDir: String

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

@@ -73,7 +73,7 @@ class MobileNet_ssd_hand: Net{
   }
   
   var preprocessKernel: CusomKernel
-  let dim = [1, 300, 300, 3]
+  let dim: (n: Int, h: Int, w: Int, c: Int) = (n: 1, h: 300, w: 300, c: 3)
   let modelPath: String
   let paramPath: String
   let modelDir: String

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

@@ -20,12 +20,17 @@ import Foundation
 import paddle_mobile
 import MetalPerformanceShaders
 
+class ScaleKernel: CusomKernel {
+  init(device: MTLDevice, shape: Shape) {
+    super.init(device: device, inFunctionName: "scale", outputDim: shape, usePaddleMobileLib: false)
+  }
+}
 
 protocol Net {
   var program: Program? { get set }
   var executor: Executor<Float32>? { get set }
   var except: Int { get }
-  var dim: [Int] { get }
+  var dim: (n: Int, h: Int, w: Int, c: Int) { get }
   var modelPath: String { get }
   var paramPath: String { get }
   var modelDir: String { get }
@@ -56,7 +61,7 @@ extension Net {
     guard let inExecutor = executor else {
       fatalError(" 请先 load ")
     }
-    try inExecutor.predict(input: inTexture, dim: dim, completionHandle: { (result) in
+    try inExecutor.predict(input: inTexture, dim: [dim.n, dim.h, dim.w, dim.c], completionHandle: { (result) in
       
       var resultArr:[Float32] = []
       resultArr = self.fetchResult(paddleMobileRes: result)
@@ -73,7 +78,7 @@ extension Net {
   
   func getTexture(image: CGImage, getTexture: @escaping (MTLTexture) -> Void) {
     let texture = try? MetalHelper.shared.textureLoader.newTexture(cgImage: image, options: [:]) ?! " texture loader error"
-    MetalHelper.scaleTexture(queue: MetalHelper.shared.queue, input: texture!, size: (224, 224)) { (resTexture) in
+    MetalHelper.scaleTexture(queue: MetalHelper.shared.queue, input: texture!, size: (dim.w, dim.h)) { (resTexture) in
       getTexture(resTexture)
     }
   }

metal/paddle-mobile-demo/paddle-mobile-demo/Net/PreProcessKernel.metal

-//
-//  PreProcessKernel.metal
-//  paddle-mobile-demo
-//
-//  Created by liuRuiLong on 2018/7/20.
-//  Copyright © 2018年 orange. All rights reserved.
-//
+/* 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;
@@ -68,10 +74,7 @@ kernel void mobilenet_ssd_preprocess_half(
 }
 
 
-kernel void genet_preprocess(
-                                     texture2d<float, access::read> inTexture [[texture(0)]],
-                                     texture2d<float, access::write> outTexture [[texture(1)]],
-                                     uint2 gid [[thread_position_in_grid]])
+kernel void genet_preprocess(texture2d<float, access::read> inTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(1)]], uint2 gid [[thread_position_in_grid]])
 {
   if (gid.x >= outTexture.get_width() ||
       gid.y >= outTexture.get_height()) {
@@ -82,3 +85,22 @@ kernel void genet_preprocess(
   outTexture.write(float4(inColor.z, inColor.y, inColor.x, 0.0f), gid);
 }
 
+kernel void scale(texture2d<float, access::sample> inTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(1)]], uint2 gid [[thread_position_in_grid]]) {
+  if (gid.x >= outTexture.get_width() ||
+      gid.y >= outTexture.get_height()) return;
+  float w_stride = inTexture.get_width() / outTexture.get_width();
+  float h_stride = inTexture.get_height() / outTexture.get_height();
+  constexpr sampler sample(coord::pixel, filter::nearest, address::clamp_to_zero);
+  float4 input = inTexture.sample(sample, float2(gid.x * w_stride,    gid.y * h_stride), 0);
+  outTexture.write(input, gid);
+}
+
+
+
+
+
+
+
+
+
+

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

@@ -20,6 +20,7 @@ import MetalPerformanceShaders
 let threadSupport = [1]
 
 let modelHelperMap: [SupportModel : Net] = [.mobilenet_ssd : MobileNet_ssd_hand.init(), .genet : Genet.init()]
+//, .genet : Genet.init()
 //let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()]
 
 enum SupportModel: String{
@@ -28,11 +29,10 @@ enum SupportModel: String{
   case genet          = "enet"
   static func supportedModels() -> [SupportModel] {
     //.mobilenet,
-    return [.mobilenet_ssd, .genet]
+    return [.mobilenet_ssd ,.genet]
   }
 }
 
-
 class ViewController: UIViewController {
   @IBOutlet weak var resultTextView: UITextView!
   @IBOutlet weak var selectImageView: UIImageView!
@@ -79,7 +79,7 @@ class ViewController: UIViewController {
       return
     }
     do {
-      let max = 1
+      let max = 50
       let startDate = Date.init()
       for i in 0..<max {
         try net.predict(inTexture: inTexture) { [weak self] (result) in

metal/paddle-mobile/paddle-mobile.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/paddle-mobile.xcscheme

@@ -33,7 +33,7 @@
       </AdditionalOptions>
    </TestAction>
    <LaunchAction
-      buildConfiguration = "Debug"
+      buildConfiguration = "Release"
       selectedDebuggerIdentifier = "Xcode.DebuggerFoundation.Debugger.LLDB"
       selectedLauncherIdentifier = "Xcode.DebuggerFoundation.Launcher.LLDB"
       launchStyle = "0"

metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift

@@ -210,7 +210,7 @@ extension MTLDevice {
 }
 
 extension MTLComputeCommandEncoder {
-  func dispatch(computePipline: MTLComputePipelineState, outTexture: MTLTexture) {
+  public func dispatch(computePipline: MTLComputePipelineState, outTexture: MTLTexture) {
     let slices = (outTexture.arrayLength * 4 + 3)/4
     
     let width = computePipline.threadExecutionWidth

metal/paddle-mobile/paddle-mobile/Operators/Kernels/Base/Kernel.swift

@@ -57,7 +57,7 @@ open class CusomKernel: Kernel {
             channel = inChannel
         }
     }
-    let outputTexture: MTLTexture
+    public let outputTexture: MTLTexture
     public init(device: MTLDevice, inFunctionName: String, outputDim: Shape, usePaddleMobileLib: Bool = false) {
         let textureDesc = MTLTextureDescriptor.init()
         textureDesc.textureType = .type2D
@@ -72,7 +72,7 @@ open class CusomKernel: Kernel {
         super.init(device: device, inFunctionName: inFunctionName, usePaddleMobileLib: usePaddleMobileLib)
     }
     
-    func compute(inputTexuture: MTLTexture, commandBuffer: MTLCommandBuffer) throws {
+    public func compute(inputTexuture: MTLTexture, commandBuffer: MTLCommandBuffer) throws {
         guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
             throw PaddleMobileError.predictError(message: " encode is nil")
         }

metal/paddle-mobile/paddle-mobile/Operators/Kernels/PoolKernel.swift

@@ -28,7 +28,7 @@ class PoolKernel<P: PrecisionType>: Kernel, Computable{
   
   required init(device: MTLDevice, param: PoolParam<P>) {
     super.init(device: device, inFunctionName: "pool")
-    param.output.initTexture(device: device)
+    param.output.initTexture(device: device, inTranspose: param.input.transpose)
   }
   
   func compute(commandBuffer: MTLCommandBuffer, param: PoolParam<P>) throws {