modify 2d dim texture

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

@@ -8,6 +8,7 @@
 
 /* Begin PBXBuildFile section */
 		30D0ED21F392CFA3885B1002 /* Pods_paddle_mobile_demo.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 18896810981724F8A0FED62A /* Pods_paddle_mobile_demo.framework */; };
+		FC013928210204A3008100E3 /* PreProcessKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC013927210204A3008100E3 /* PreProcessKernel.metal */; };
 		FC039B8220E11C550081E9F8 /* AppDelegate.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC039B8120E11C550081E9F8 /* AppDelegate.swift */; };
 		FC039B8420E11C550081E9F8 /* ViewController.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC039B8320E11C550081E9F8 /* ViewController.swift */; };
 		FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8520E11C550081E9F8 /* Main.storyboard */; };
@@ -238,6 +239,7 @@
 		081C9CF10DB06C58B8B6B039 /* Pods-paddle-mobile-demo.release.xcconfig */ = {isa = PBXFileReference; includeInIndex = 1; lastKnownFileType = text.xcconfig; name = "Pods-paddle-mobile-demo.release.xcconfig"; path = "../Pods/Target Support Files/Pods-paddle-mobile-demo/Pods-paddle-mobile-demo.release.xcconfig"; sourceTree = "<group>"; };
 		18896810981724F8A0FED62A /* Pods_paddle_mobile_demo.framework */ = {isa = PBXFileReference; explicitFileType = wrapper.framework; includeInIndex = 0; path = Pods_paddle_mobile_demo.framework; sourceTree = BUILT_PRODUCTS_DIR; };
 		878829884E1A14D7044721D5 /* Pods-paddle-mobile-demo.debug.xcconfig */ = {isa = PBXFileReference; includeInIndex = 1; lastKnownFileType = text.xcconfig; name = "Pods-paddle-mobile-demo.debug.xcconfig"; path = "../Pods/Target Support Files/Pods-paddle-mobile-demo/Pods-paddle-mobile-demo.debug.xcconfig"; sourceTree = "<group>"; };
+		FC013927210204A3008100E3 /* PreProcessKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = PreProcessKernel.metal; sourceTree = "<group>"; };
 		FC039B7E20E11C550081E9F8 /* paddle-mobile-demo.app */ = {isa = PBXFileReference; explicitFileType = wrapper.application; includeInIndex = 0; path = "paddle-mobile-demo.app"; sourceTree = BUILT_PRODUCTS_DIR; };
 		FC039B8120E11C550081E9F8 /* AppDelegate.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = AppDelegate.swift; sourceTree = "<group>"; };
 		FC039B8320E11C550081E9F8 /* ViewController.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ViewController.swift; sourceTree = "<group>"; };
@@ -506,6 +508,7 @@
 				FC0E2C2020EDC03B009C1FAC /* models */,
 				FC0E2C1D20EDC030009C1FAC /* images */,
 				FC039B8120E11C550081E9F8 /* AppDelegate.swift */,
+				FC013927210204A3008100E3 /* PreProcessKernel.metal */,
 				FC039B8320E11C550081E9F8 /* ViewController.swift */,
 				FC039B8520E11C550081E9F8 /* Main.storyboard */,
 				FC039B8820E11C560081E9F8 /* Assets.xcassets */,
@@ -1067,6 +1070,7 @@
 			buildActionMask = 2147483647;
 			files = (
 				FC039B8420E11C550081E9F8 /* ViewController.swift in Sources */,
+				FC013928210204A3008100E3 /* PreProcessKernel.metal in Sources */,
 				FC039B8220E11C550081E9F8 /* AppDelegate.swift in Sources */,
 			);
 			runOnlyForDeploymentPostprocessing = 0;

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

+//
+//  PreProcessKernel.metal
+//  paddle-mobile-demo
+//
+//  Created by liuRuiLong on 2018/7/20.
+//  Copyright © 2018年 orange. All rights reserved.
+//
+
+#include <metal_stdlib>
+using namespace metal;
+
+
+kernel void 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()) {
+        return;
+    }
+    // Subtract mean values, scale by 0.017, convert to BGR.
+    
+    const auto means = float4(103.94f, 116.78f, 123.68f, 0.0f);
+    const float4 inColor = (float4(inTexture.read(gid)) * 255.0f - means) * 0.017f;
+    outTexture.write(float4(inColor.x, inColor.y, inColor.z, 0.0f), gid);
+}
+

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

@@ -17,12 +17,14 @@ import MetalKit
 import paddle_mobile
 import MetalPerformanceShaders
 
-
-
-func Test<T>() -> T? {
-    return nil
+class PreProccess: CusomKernel {
+    init(device: MTLDevice) {
+        let s = CusomKernel.Shape.init(inWidth: 224, inHeight: 224, inChannel: 3)
+        super.init(device: device, inFunctionName: "preprocess", outputDim: s, usePaddleMobileLib: false)
+    }
 }
 
+
 class ViewController: UIViewController {
     let device: MTLDevice! = MTLCreateSystemDefaultDevice()
     var textureLoader: MTKTextureLoader!
@@ -63,6 +65,8 @@ class ViewController: UIViewController {
         guard let inTexture = texture else {
             fatalError(" texture is nil !")
         }
+        
+        
        
         scaleTexture(queue: queue!, input: inTexture) { (inputTexture) in
             let loader = Loader<Float32>.init()
@@ -71,7 +75,8 @@ class ViewController: UIViewController {
                 let paraPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null"
                 let program = try loader.load(device: self.device, modelPath: modelPath, paraPath: paraPath)
                 let executor = try Executor<Float32>.init(inDevice: self.device, inQueue: queue!, inProgram: program)
-                let output = try executor.predict(input: inputTexture, expect: [1, 224, 224, 3])
+                let preprocessKernel = PreProccess.init(device: self.device)
+                let output = try executor.predict(input: inputTexture, expect: [1, 224, 224, 3], preProcessKernle: preprocessKernel)
                 //            print(output)
             } catch let error {
                 print(error)

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

@@ -101,7 +101,7 @@ public extension MTLTexture {
                 getBytes(bytes, bytesPerRow: bytesPerRow, bytesPerImage: bytesPerImage, from: region, mipmapLevel: 0, slice: i)
                 let p = bytes.assumingMemoryBound(to: T.self)
                 str += "2d array count : \(width * height * depth * 4) \n"
-                if stridable {
+                if stridable && width * height * depth * 4 > 100 {
                     for j in stride(from: 0, to: width * height * depth * 4 , by: width * height * depth * 4 / 100){
                         str += " index \(j): \(p[j])"
                     }

metal/paddle-mobile/paddle-mobile/Executor.swift

@@ -79,14 +79,27 @@ public class Executor<P: PrecisionType> {
         }
     }
     
-    public func predict(input: MTLTexture, expect: [Int]) throws -> ResultHolder<P> {
-        let beforeDate = Date.init()
-        let inputTexture = InputTexture.init(inMTLTexture: input, inExpectDim: Dim.init(inDim: expect))
-        program.scope.setInput(input: inputTexture)
+    public func predict(input: MTLTexture, expect: [Int], preProcessKernle: CusomKernel? = nil) throws -> ResultHolder<P> {
         guard let buffer = queue.makeCommandBuffer() else {
             throw PaddleMobileError.predictError(message: "CommandBuffer is nil")
         }
+        let resInput: MTLTexture
         
+        if let inPre = preProcessKernle {
+            do {
+                try inPre.compute(inputTexuture: input, commandBuffer: buffer)
+                resInput = inPre.outputTexture
+            } catch let error {
+                throw error
+            }
+        } else {
+            resInput = input
+        }
+        
+        let beforeDate = Date.init()
+        let inputTexture = InputTexture.init(inMTLTexture: resInput, inExpectDim: Dim.init(inDim: expect))
+        program.scope.setInput(input: inputTexture)
+ 
         for op in ops {
             do {
                 try op.run(device: device, buffer: buffer)
@@ -100,7 +113,6 @@ public class Executor<P: PrecisionType> {
                 op.delogOutput()
             }
             
-
             let afterDate = Date.init()
             print(" encoder end ! time: \(afterDate.timeIntervalSince(beforeDate))")
         }

metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift

@@ -107,16 +107,17 @@ class ConvAddBatchNormReluOp<P: PrecisionType>: Operator<ConvAddBatchNormReluKer
     }
     
     func delogOutput() {
-//        let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: true)
-//        para.filter.logDataPointer(header: "filter data pointer: ")
         
+//        let _: P? = para.input.metalTexture.logDesc(header: "conv add batchnorm relu input: ", stridable: false)
+//        para.filter.logDataPointer(header: "filter data pointer: ")
 //        print("filter: \(para.filter)")
-//        print("biase: \(para.bias)")
-//        print("padding: \(para.paddings)")
-//        print("stride: \(para.stride)")
         
-//        let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false)
-//        let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false)
+        print("biase: \(para.y)")
+        print("padding: \(para.paddings)")
+        print("stride: \(para.stride)")
+        
+        let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: true)
+        let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: true)
 //        let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false)
     }
 }

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

@@ -33,7 +33,7 @@ class ConvAddKernel<P: PrecisionType>: Kernel, Computable {
             throw PaddleMobileError.predictError(message: " encode is nil")
         }
         
-        print("Conv Add compute")
+        print("Conv add compute")
         encoder.setTexture(param.input.metalTexture, index: 0)
         encoder.setTexture(param.output.metalTexture, index: 1)
         encoder.setBytes(&metalParam, length: MemoryLayout<MetalConvParam>.size, index: 0)

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

@@ -33,7 +33,6 @@ class ConvKernel<P: PrecisionType>: Kernel, Computable {
         let offsetZ = 0.0
         
         metalParam = MetalConvParam.init(offsetX: Int16(offsetX), offsetY: Int16(offsetY), offsetZ: Int16(offsetZ), strideX: UInt16(param.stride[0]), strideY: UInt16(param.stride[1]), paddedZ: UInt16(param.input.metalTexture.arrayLength * 4 - param.input.dim[3]))
-        
     }
     
     func compute(commandBuffer: MTLCommandBuffer, param: ConvParam<P>) throws {

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

@@ -27,12 +27,50 @@ protocol KernelProtocol {
    
 }
 
-class Kernel {
+open class Kernel {
     let pipline: MTLComputePipelineState
     let functionName: String
-    init(device: MTLDevice, inFunctionName: String) {
-        pipline = device.pipeLine(funcName: inFunctionName)
+    public init(device: MTLDevice, inFunctionName: String, usePaddleMobileLib: Bool = true) {
+        pipline = device.pipeLine(funcName: inFunctionName, inPaddleMobileLib: usePaddleMobileLib)
         functionName = inFunctionName
     }
 }
 
+open class CusomKernel: Kernel {
+    public struct Shape {
+        public let width: Int
+        public let height: Int
+        public let channel: Int
+        public init(inWidth: Int, inHeight: Int, inChannel: Int){
+            width = inWidth
+            height = inHeight
+            channel = inChannel
+        }
+    }
+    let outputTexture: MTLTexture
+    public init(device: MTLDevice, inFunctionName: String, outputDim: Shape, usePaddleMobileLib: Bool = false) {
+        let textureDesc = MTLTextureDescriptor.init()
+        textureDesc.textureType = .type2D
+        textureDesc.width = outputDim.width
+        textureDesc.height = outputDim.height
+        textureDesc.depth = (outputDim.channel + 3) / 4
+        textureDesc.pixelFormat = .rgba32Float
+        textureDesc.usage = [.shaderRead, .shaderWrite]
+        textureDesc.storageMode = .shared
+        outputTexture = device.makeTexture(descriptor: textureDesc) ?! " make texture error "
+
+        super.init(device: device, inFunctionName: inFunctionName, usePaddleMobileLib: usePaddleMobileLib)
+    }
+    
+    func compute(inputTexuture: MTLTexture, commandBuffer: MTLCommandBuffer) throws {
+        guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
+            throw PaddleMobileError.predictError(message: " encode is nil")
+        }
+        encoder.setTexture(inputTexuture, index: 0)
+        encoder.setTexture(outputTexture, index: 1)
+        encoder.dispatch(computePipline: pipline, outTexture: outputTexture)
+        encoder.endEncoding()
+    }
+    
+}
+

metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift

@@ -58,11 +58,11 @@ class PoolOp<P: PrecisionType>: Operator<PoolKernel<P>, PoolParam<P>>, Runable,
     
     func delogOutput() {
         print("pool2d delog")
-        let _: P? = para.input.metalTexture.logDesc(header: "pool2d input: ", stridable: false)
+        let _: P? = para.input.metalTexture.logDesc(header: "pool2d input: ", stridable: true)
         print(para.ksize)
         print(para.stride)
         print(para.padding)
         print(para.poolType)
-        let _: P? = para.output.metalTexture.logDesc(header: "pool2d output: ", stridable: false)
+        let _: P? = para.output.metalTexture.logDesc(header: "pool2d output: ", stridable: true)
     }
 }

metal/paddle-mobile/paddle-mobile/framework/Texture.swift

@@ -59,10 +59,11 @@ public class Texture<P: PrecisionType>: Tensorial {
             tmpTextureDes.arrayLength = (inDim[3] * inDim[0] + 3)/4
             tmpTextureDes.textureType = .type2DArray
         } else if inDim.cout() == 2 {
-            tmpTextureDes.height = inDim[0]
-            tmpTextureDes.width = inDim[1]
+            tmpTextureDes.height = 1
+            tmpTextureDes.width = 1
             tmpTextureDes.depth = 1
-            tmpTextureDes.textureType = .type2D
+            tmpTextureDes.arrayLength = (inDim[0] * inDim[1] + 3)/4
+            tmpTextureDes.textureType = .type2DArray
         } else {
             fatalError(" not suuprt ")
         }