From bbced1e4a77fdc0913137d460f664c9cfe4309cd Mon Sep 17 00:00:00 2001 From: liuruilong Date: Mon, 23 Jul 2018 19:29:51 +0800 Subject: [PATCH] modify 2d dim texture --- .../project.pbxproj | 4 ++ .../paddle-mobile-demo/PreProcessKernel.metal | 28 ++++++++++++ .../paddle-mobile-demo/ViewController.swift | 15 ++++--- .../paddle-mobile/Common/MetalExtension.swift | 2 +- .../paddle-mobile/Executor.swift | 22 +++++++--- .../Operators/ConvAddBatchNormReluOp.swift | 15 ++++--- .../Operators/Kernels/ConvAddKernel.swift | 2 +- .../Operators/Kernels/ConvKernel.swift | 1 - .../Operators/Kernels/Kernel.swift | 44 +++++++++++++++++-- .../paddle-mobile/Operators/PoolOp.swift | 4 +- .../paddle-mobile/framework/Texture.swift | 7 +-- 11 files changed, 116 insertions(+), 28 deletions(-) create mode 100644 metal/paddle-mobile-demo/paddle-mobile-demo/PreProcessKernel.metal diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj b/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj index 1089a98fc5..e5a0e3bc21 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj +++ b/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 = ""; }; 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 = ""; }; + FC013927210204A3008100E3 /* PreProcessKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = PreProcessKernel.metal; sourceTree = ""; }; 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 = ""; }; FC039B8320E11C550081E9F8 /* ViewController.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ViewController.swift; sourceTree = ""; }; @@ -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; diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/PreProcessKernel.metal b/metal/paddle-mobile-demo/paddle-mobile-demo/PreProcessKernel.metal new file mode 100644 index 0000000000..f32ab30d6b --- /dev/null +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/PreProcessKernel.metal @@ -0,0 +1,28 @@ +// +// PreProcessKernel.metal +// paddle-mobile-demo +// +// Created by liuRuiLong on 2018/7/20. +// Copyright © 2018年 orange. All rights reserved. +// + +#include +using namespace metal; + + +kernel void preprocess( + texture2d inTexture [[texture(0)]], + texture2d 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); +} + diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift index c84d8a7e5a..74ad1c7a51 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift @@ -17,12 +17,14 @@ import MetalKit import paddle_mobile import MetalPerformanceShaders - - -func Test() -> 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.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.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) diff --git a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift index 07127defa6..5ae0d66470 100644 --- a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift +++ b/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])" } diff --git a/metal/paddle-mobile/paddle-mobile/Executor.swift b/metal/paddle-mobile/paddle-mobile/Executor.swift index ec21b51487..00084a63cb 100644 --- a/metal/paddle-mobile/paddle-mobile/Executor.swift +++ b/metal/paddle-mobile/paddle-mobile/Executor.swift @@ -79,14 +79,27 @@ public class Executor { } } - public func predict(input: MTLTexture, expect: [Int]) throws -> ResultHolder

{ - 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

{ 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 { op.delogOutput() } - let afterDate = Date.init() print(" encoder end ! time: \(afterDate.timeIntervalSince(beforeDate))") } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift index 5c78b03f9f..0968888c8b 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift @@ -107,16 +107,17 @@ class ConvAddBatchNormReluOp: Operator: 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.size, index: 0) diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvKernel.swift index 65a62121a6..a61fd4f048 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvKernel.swift @@ -33,7 +33,6 @@ class ConvKernel: 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

) throws { diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernel.swift index f3790ff695..a2d93ac25a 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernel.swift +++ b/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() + } + +} + diff --git a/metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift index fde35e3a83..07676defe7 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift @@ -58,11 +58,11 @@ class PoolOp: Operator, PoolParam

>, 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) } } diff --git a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift index 9370092e1d..eae373a141 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift @@ -59,10 +59,11 @@ public class Texture: 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 ") } -- GitLab