From 0ca8fdeaef7abc49d2c5bab7648c2212091e0872 Mon Sep 17 00:00:00 2001 From: liuruilong Date: Sat, 29 Dec 2018 12:12:05 +0800 Subject: [PATCH] support super resulotion fix metal popool op bug --- .../project.pbxproj | 8 ++ .../Net/BufferToTexture.metal | 35 +++++++++ .../paddle-mobile.xcodeproj/project.pbxproj | 4 + .../paddle-mobile/CustomNet/Genet.swift | 2 +- .../paddle-mobile/CustomNet/MobileNet.swift | 2 +- .../CustomNet/MobileNetCombined.swift | 8 +- .../CustomNet/MobileNetSSD.swift | 2 +- .../CustomNet/MobilenetSSD_AR.swift | 2 +- .../CustomNet/SuperResolutionNet.swift | 8 ++ .../paddle-mobile/CustomNet/YoloNet.swift | 2 +- .../Operators/Kernels/Base/Kernel.swift | 61 +++++++++++++--- .../Operators/Kernels/PoolKernel.swift | 2 +- .../Operators/Kernels/metal/Macro.metal | 2 +- .../Kernels/metal/PoolKernel.inc.metal | 44 +++++++++++ .../Operators/Kernels/metal/PoolKernel.metal | 73 ++----------------- .../paddle-mobile/paddle-mobile/Runner.swift | 30 +++++++- .../paddle-mobile/framework/Executor.swift | 15 ++-- .../paddle-mobile/framework/Texture.swift | 2 - src/framework/executor.cpp | 2 +- 19 files changed, 204 insertions(+), 100 deletions(-) create mode 100644 metal/paddle-mobile-demo/paddle-mobile-demo/Net/BufferToTexture.metal create mode 100644 metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PoolKernel.inc.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 4de49759d0..bee3fa3c18 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj +++ b/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj @@ -29,6 +29,8 @@ FC9797C221D608E000F2FD90 /* mobilenet_model in Resources */ = {isa = PBXBuildFile; fileRef = FC9797C021D608DF00F2FD90 /* mobilenet_model */; }; FC9797C321D608E000F2FD90 /* mobilenet_params in Resources */ = {isa = PBXBuildFile; fileRef = FC9797C121D608DF00F2FD90 /* mobilenet_params */; }; FC9797C721D609FB00F2FD90 /* synset.txt in Resources */ = {isa = PBXBuildFile; fileRef = FC9797C621D609FB00F2FD90 /* synset.txt */; }; + FC9797CF21D6506F00F2FD90 /* mingren.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FC9797CE21D6506F00F2FD90 /* mingren.jpg */; }; + FC9797D121D6616600F2FD90 /* BufferToTexture.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC9797D021D6616600F2FD90 /* BufferToTexture.metal */; }; FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */ = {isa = PBXBuildFile; fileRef = FCBCCC542122EF5400D94F7E /* MetalHelper.swift */; }; 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, ); }; }; @@ -78,6 +80,8 @@ FC9797C021D608DF00F2FD90 /* mobilenet_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = mobilenet_model; sourceTree = ""; }; FC9797C121D608DF00F2FD90 /* mobilenet_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = mobilenet_params; sourceTree = ""; }; FC9797C621D609FB00F2FD90 /* synset.txt */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = synset.txt; sourceTree = ""; }; + FC9797CE21D6506F00F2FD90 /* mingren.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = mingren.jpg; sourceTree = ""; }; + FC9797D021D6616600F2FD90 /* BufferToTexture.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = BufferToTexture.metal; sourceTree = ""; }; FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = ""; }; 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 = ""; }; @@ -158,6 +162,7 @@ FC203FA821CBFDBA00B37166 /* images */ = { isa = PBXGroup; children = ( + FC9797CE21D6506F00F2FD90 /* mingren.jpg */, FC9797BD21D6045B00F2FD90 /* banana.jpeg */, FC203FA921CBFDBA00B37166 /* test.jpg */, ); @@ -226,6 +231,7 @@ children = ( FC013927210204A3008100E3 /* PreProcessKernel.metal */, FCBCCC542122EF5400D94F7E /* MetalHelper.swift */, + FC9797D021D6616600F2FD90 /* BufferToTexture.metal */, ); path = Net; sourceTree = ""; @@ -303,6 +309,7 @@ buildActionMask = 2147483647; files = ( FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */, + FC9797CF21D6506F00F2FD90 /* mingren.jpg in Resources */, FC704C2221D237FC00F98BAB /* combined_mobilenet_params in Resources */, FC704C1921D2375300F98BAB /* super_params in Resources */, FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */, @@ -373,6 +380,7 @@ FC803BCD214D27930094B8E5 /* FPSCounter.swift in Sources */, C2E67E5E21524E460013F575 /* LoadPointerViewController.m in Sources */, FC039B8220E11C550081E9F8 /* AppDelegate.swift in Sources */, + FC9797D121D6616600F2FD90 /* BufferToTexture.metal in Sources */, ); runOnlyForDeploymentPostprocessing = 0; }; diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/Net/BufferToTexture.metal b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/BufferToTexture.metal new file mode 100644 index 0000000000..97969100c0 --- /dev/null +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/Net/BufferToTexture.metal @@ -0,0 +1,35 @@ +// +// RGBToYCrCb_Y.metal +// paddle-mobile-demo +// +// Created by liuRuiLong on 2018/12/28. +// Copyright © 2018 orange. All rights reserved. +// + +#include +using namespace metal; + +kernel void buffer_to_texture_kernel( const device float *input [[buffer(0)]], +texture2d outTexture [[texture(0)]], +uint2 gid [[thread_position_in_grid]]){ + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height()) { + return; + } + + float y = input[outTexture.get_width() * gid.y + gid.x]; + outTexture.write(float4(y, 0.0f, 0.0f, 0.0f), gid); +} + +kernel void buffer_to_texture_kernel_half( const device float *input [[buffer(0)]], + texture2d outTexture [[texture(0)]], + uint2 gid [[thread_position_in_grid]]){ + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height()) { + return; + } + + float y = input[outTexture.get_width() * gid.y + gid.x]; + outTexture.write(half4(y, 0.0f, 0.0f, 0.0f), gid); +} + diff --git a/metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj b/metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj index cf89c365da..41e30cd0f7 100644 --- a/metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj +++ b/metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj @@ -112,6 +112,7 @@ 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 */; }; @@ -254,6 +255,7 @@ FCBCCC6C2123073A00D94F7E /* BoxcoderKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = BoxcoderKernel.swift; sourceTree = ""; }; FCBCCC6E2123097100D94F7E /* MulticlassNMSOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MulticlassNMSOp.swift; sourceTree = ""; }; FCBCCC70212309A700D94F7E /* MulticlassNMSKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MulticlassNMSKernel.swift; sourceTree = ""; }; + FCCED5E021D71FC000BE8D5F /* PoolKernel.inc.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = PoolKernel.inc.metal; sourceTree = ""; }; FCD04E6520F314C50007374F /* PoolOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PoolOp.swift; sourceTree = ""; }; FCD04E6720F315020007374F /* PoolKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PoolKernel.swift; sourceTree = ""; }; FCD04E6920F319EC0007374F /* SoftmaxOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = SoftmaxOp.swift; sourceTree = ""; }; @@ -535,6 +537,7 @@ FCA67CD82138287B00BD58AA /* ConvBNReluKernel.metal */, FC0226552138F33800F395E2 /* TransposeKernel.metal */, 4AA1EAAD214F5FD900D0F791 /* TransposeKernel.inc.metal */, + FCCED5E021D71FC000BE8D5F /* PoolKernel.inc.metal */, FC0226572138F38D00F395E2 /* PoolKernel.metal */, FC803BC2214CB79C0094B8E5 /* ConvAddPreluKernel.metal */, FC803BC4214CB8F00094B8E5 /* ConvAddPrelu.inc.metal */, @@ -707,6 +710,7 @@ FCA3A1632132A4AC00084FE5 /* ReshapeKernel.metal in Sources */, FCBCCC592122F42700D94F7E /* ConvBNReluOp.swift in Sources */, FC039BA920E11CBC0081E9F8 /* ConvOp.swift in Sources */, + FCCED5E121D71FC000BE8D5F /* PoolKernel.inc.metal in Sources */, FC9D038420E23B01000F735A /* Texture.swift in Sources */, FCE3A1B32153E91900C37CDE /* ElementwiseAddPreluKernel.metal in Sources */, 4AA1EAA2214912CD00D0F791 /* FlattenKernel.swift in Sources */, diff --git a/metal/paddle-mobile/paddle-mobile/CustomNet/Genet.swift b/metal/paddle-mobile/paddle-mobile/CustomNet/Genet.swift index 8cb9f680e5..6babcc744d 100644 --- a/metal/paddle-mobile/paddle-mobile/CustomNet/Genet.swift +++ b/metal/paddle-mobile/paddle-mobile/CustomNet/Genet.swift @@ -41,7 +41,7 @@ public class Genet: Net { class GenetPreProccess: CusomKernel { init(device: MTLDevice) { - let s = CusomKernel.Shape.init(inWidth: 128, inHeight: 128, inChannel: 3) + let s = Shape.init(inWidth: 128, inHeight: 128, inChannel: 3) super.init(device: device, inFunctionName: "genet_preprocess", outputDim: s, usePaddleMobileLib: false) } } diff --git a/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNet.swift b/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNet.swift index 9acc8f602c..db79480008 100644 --- a/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNet.swift +++ b/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNet.swift @@ -18,7 +18,7 @@ public class MobileNet: Net{ class MobilenetPreProccess: CusomKernel { init(device: MTLDevice) { - let s = CusomKernel.Shape.init(inWidth: 224, inHeight: 224, inChannel: 3) + let s = Shape.init(inWidth: 224, inHeight: 224, inChannel: 3) super.init(device: device, inFunctionName: "mobilenet_preprocess", outputDim: s, usePaddleMobileLib: false) } } diff --git a/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNetCombined.swift b/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNetCombined.swift index 8c89ec8e58..f0f871c960 100644 --- a/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNetCombined.swift +++ b/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNetCombined.swift @@ -17,11 +17,11 @@ 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" modelDir = "" - //preprocessKernel = GenetPreProccess.init(device: device) - inputDim_ = Dim.init(inDim: [1, 416, 416, 3]) + inputDim_ = Dim.init(inDim: [1, 224, 224, 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) means = [0, 0, 0] scale = 1 @@ -29,9 +29,7 @@ public class MobileNetCombined: Net { modelPath = "" paramPath = "" modelDir = "" - //preprocessKernel = GenetPreProccess.init(device: device) - inputDim_ = Dim.init(inDim: [1, 416, 416, 3]) - + inputDim_ = Dim.init(inDim: [1, 224, 224, 3]) } // class GenetPreProccess: CusomKernel { diff --git a/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNetSSD.swift b/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNetSSD.swift index e20b88068b..93e5b469e0 100644 --- a/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNetSSD.swift +++ b/metal/paddle-mobile/paddle-mobile/CustomNet/MobileNetSSD.swift @@ -41,7 +41,7 @@ public class MobileNet_ssd_hand: Net{ class MobilenetssdPreProccess: CusomKernel { init(device: MTLDevice) { - let s = CusomKernel.Shape.init(inWidth: 300, inHeight: 300, inChannel: 3) + let s = Shape.init(inWidth: 300, inHeight: 300, inChannel: 3) super.init(device: device, inFunctionName: "mobilenet_ssd_preprocess", outputDim: s, usePaddleMobileLib: false) } } diff --git a/metal/paddle-mobile/paddle-mobile/CustomNet/MobilenetSSD_AR.swift b/metal/paddle-mobile/paddle-mobile/CustomNet/MobilenetSSD_AR.swift index 2b8e1d1729..56af7339cb 100644 --- a/metal/paddle-mobile/paddle-mobile/CustomNet/MobilenetSSD_AR.swift +++ b/metal/paddle-mobile/paddle-mobile/CustomNet/MobilenetSSD_AR.swift @@ -41,7 +41,7 @@ public class MobileNet_ssd_AR: Net{ class MobilenetssdPreProccess: CusomKernel { init(device: MTLDevice) { - let s = CusomKernel.Shape.init(inWidth: 160, inHeight: 160, inChannel: 3) + let s = Shape.init(inWidth: 160, inHeight: 160, inChannel: 3) super.init(device: device, inFunctionName: "mobilent_ar_preprocess", outputDim: s, usePaddleMobileLib: false) } } diff --git a/metal/paddle-mobile/paddle-mobile/CustomNet/SuperResolutionNet.swift b/metal/paddle-mobile/paddle-mobile/CustomNet/SuperResolutionNet.swift index 134aa3b9d8..9df8afb977 100644 --- a/metal/paddle-mobile/paddle-mobile/CustomNet/SuperResolutionNet.swift +++ b/metal/paddle-mobile/paddle-mobile/CustomNet/SuperResolutionNet.swift @@ -14,6 +14,14 @@ import Foundation + +class SuperResolutionPreProccess: CusomKernel { + init(device: MTLDevice) { + let s = Shape.init(inWidth: 224, inHeight: 224, inChannel: 3) + super.init(device: device, inFunctionName: "super_resolution_preprocess", outputDim: s, usePaddleMobileLib: false) + } +} + public class SuperResolutionNet: Net{ override public func resultStr(res: ResultHolder) -> String { return "未实现" diff --git a/metal/paddle-mobile/paddle-mobile/CustomNet/YoloNet.swift b/metal/paddle-mobile/paddle-mobile/CustomNet/YoloNet.swift index 105654ad76..0cb1cb7755 100644 --- a/metal/paddle-mobile/paddle-mobile/CustomNet/YoloNet.swift +++ b/metal/paddle-mobile/paddle-mobile/CustomNet/YoloNet.swift @@ -18,7 +18,7 @@ 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" modelDir = "" - //preprocessKernel = GenetPreProccess.init(device: device) +// preprocessKernel = GenetPreProccess.init(device: device) inputDim_ = Dim.init(inDim: [1, 416, 416, 3]) } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Base/Kernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Base/Kernel.swift index f58358761f..06a4caf2ce 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Base/Kernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Base/Kernel.swift @@ -46,17 +46,60 @@ open class Kernel { } } -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 +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 + } +} + +open class BufferToTextureKernel: Kernel { + public let outputTexture: MTLTexture + + public init(device: MTLDevice, 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 + + if computePrecision == .Float16 { + textureDesc.pixelFormat = .rgba16Float + } else if computePrecision == .Float32 { + textureDesc.pixelFormat = .rgba32Float + } else { + fatalError() } + + textureDesc.usage = [.shaderRead, .shaderWrite] + textureDesc.storageMode = .shared + outputTexture = device.makeTexture(descriptor: textureDesc) ?! " make texture error " + if computePrecision == .Float32 { + super.init(device: device, inFunctionName: "buffer_to_texture_kernel", usePaddleMobileLib: usePaddleMobileLib) + } else { + super.init(device: device, inFunctionName: "buffer_to_texture_kernel_half", usePaddleMobileLib: usePaddleMobileLib) + } + } + + public func compute(inputBuffer: MTLBuffer , commandBuffer: MTLCommandBuffer) throws { + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { + throw PaddleMobileError.predictError(message: " encode is nil") + } + + encoder.setBuffer(inputBuffer, offset: 0, index: 0) + encoder.setTexture(outputTexture, index: 0) + encoder.dispatch(computePipline: pipline, outTexture: outputTexture) + encoder.endEncoding() } + +} + +open class CusomKernel: Kernel { + public let outputTexture: MTLTexture public init(device: MTLDevice, inFunctionName: String, outputDim: Shape, usePaddleMobileLib: Bool = false) { let textureDesc = MTLTextureDescriptor.init() diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PoolKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PoolKernel.swift index 1d66e420e2..c6c243eab1 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PoolKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PoolKernel.swift @@ -49,7 +49,7 @@ class PoolKernel: Kernel, Computable{ ) if computePrecision == .Float32 { - super.init(device: device, inFunctionName: "pool") + super.init(device: device, inFunctionName: "pool_float") } else if computePrecision == .Float16 { super.init(device: device, inFunctionName: "pool_half") } else { diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Macro.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Macro.metal index 950d7d5f05..43b9579c89 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Macro.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Macro.metal @@ -24,6 +24,6 @@ using namespace metal; #define FUNC(f, r, n, v, p) CONCAT5_(f, r, n, v, p) #define VECTOR(p, n) CONCAT2(p, n) - +#define FUNC2_(a, b) CONCAT2_(a, b) #define FUNC3_(a, b, c) CONCAT3_(a, b, c) diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PoolKernel.inc.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PoolKernel.inc.metal new file mode 100644 index 0000000000..5b0cae3e3a --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PoolKernel.inc.metal @@ -0,0 +1,44 @@ +// +// PoolKernel.inc.metal +// paddle-mobile +// +// Created by liuRuiLong on 2018/12/29. +// Copyright © 2018 orange. All rights reserved. +// + +#ifdef P + +kernel void FUNC2_(pool, P)(texture2d_array inTexture [[texture(0)]], + texture2d_array outTexture [[texture(1)]], + constant PoolParam &pm [[buffer(0)]], + uint3 gid [[thread_position_in_grid]]) { + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height() || + gid.z >= outTexture.get_array_size()) return; + int xmin = gid.x * pm.strideX - pm.paddingX; + int xmax = min(xmin + pm.ksizeX, int(inTexture.get_width())); + xmin = max(xmin, 0); + int ymin = gid.y * pm.strideX - pm.paddingX; + int ymax = min(ymin + pm.ksizeX, int(inTexture.get_height())); + ymin = max(ymin, 0); + + VECTOR(P, 4) r = 0; + if (pm.poolType == 0) { + r = inTexture.read(uint2(xmin, ymin), gid.z); + for (int x = xmin; x < xmax; x++) { + for (int y = ymin; y < ymax; y++) { + r = fmax(r, inTexture.read(uint2(x, y), gid.z)); + } + } + } else if (pm.poolType == 1) { + for (int x = xmin; x < xmax; x++) { + for (int y = ymin; y < ymax; y++) { + r += inTexture.read(uint2(x, y), gid.z); + } + } + r /= (xmax - xmin) * (ymax - ymin); + } + outTexture.write(r, gid.xy, gid.z); +} + +#endif diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PoolKernel.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PoolKernel.metal index 1f2f7240db..d17536eab6 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PoolKernel.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PoolKernel.metal @@ -13,7 +13,8 @@ limitations under the License. */ #include -#include "Common.metal" +#include "Macro.metal" + using namespace metal; struct PoolParam { @@ -26,68 +27,10 @@ struct PoolParam { int poolType; }; -kernel void pool(texture2d_array inTexture [[texture(0)]], - texture2d_array outTexture [[texture(1)]], - constant PoolParam &pm [[buffer(0)]], - uint3 gid [[thread_position_in_grid]]) { - if (gid.x >= outTexture.get_width() || - gid.y >= outTexture.get_height() || - gid.z >= outTexture.get_array_size()) return; - int xmin = gid.x * pm.strideX - pm.paddingX; - int xmax = min(xmin + pm.ksizeX, int(inTexture.get_width())); - xmin = max(xmin, 0); - int ymin = gid.y * pm.strideX - pm.paddingX; - int ymax = min(ymin + pm.ksizeX, int(inTexture.get_height())); - ymin = max(ymin, 0); - - float4 r = 0; - if (pm.poolType == 0) { - r = inTexture.read(uint2(xmin, ymin), gid.z); - for (int x = xmin; x < xmax; x++) { - for (int y = ymin; y < ymax; y++) { - r = fmax(r, inTexture.read(uint2(x, y), gid.z)); - } - } - } else if (pm.poolType == 1) { - for (int x = xmin; x < xmax; x++) { - for (int y = ymin; y < ymax; y++) { - r += inTexture.read(uint2(x, y), gid.z); - } - } - r /= pm.ksizeX * pm.ksizeY; - } - outTexture.write(r, gid.xy, gid.z); -} +#define P float +#import "PoolKernel.inc.metal" +#undef P -kernel void pool_half(texture2d_array inTexture [[texture(0)]], - texture2d_array outTexture [[texture(1)]], - constant PoolParam &pm [[buffer(0)]], - uint3 gid [[thread_position_in_grid]]) { - if (gid.x >= outTexture.get_width() || - gid.y >= outTexture.get_height() || - gid.z >= outTexture.get_array_size()) return; - int xmin = gid.x * pm.strideX - pm.paddingX; - int xmax = min(xmin + pm.ksizeX, int(inTexture.get_width())); - xmin = max(xmin, 0); - int ymin = gid.y * pm.strideX - pm.paddingX; - int ymax = min(ymin + pm.ksizeX, int(inTexture.get_height())); - ymin = max(ymin, 0); - - half4 r = 0; - if (pm.poolType == 0) { - r = inTexture.read(uint2(xmin, ymin), gid.z); - for (int x = xmin; x < xmax; x++) { - for (int y = ymin; y < ymax; y++) { - r = fmax(r, inTexture.read(uint2(x, y), gid.z)); - } - } - } else if (pm.poolType == 1) { - for (int x = xmin; x < xmax; x++) { - for (int y = ymin; y < ymax; y++) { - r += inTexture.read(uint2(x, y), gid.z); - } - } - r /= pm.ksizeX * pm.ksizeY; - } - outTexture.write(r, gid.xy, gid.z); -} +#define P half +#import "PoolKernel.inc.metal" +#undef P diff --git a/metal/paddle-mobile/paddle-mobile/Runner.swift b/metal/paddle-mobile/paddle-mobile/Runner.swift index 766e6ec9f6..0fefadd22a 100644 --- a/metal/paddle-mobile/paddle-mobile/Runner.swift +++ b/metal/paddle-mobile/paddle-mobile/Runner.swift @@ -51,7 +51,10 @@ class ScaleKernel: CusomKernel { } numel = net.inputDim.numel() - meansNumber = net.means.map { NSNumber.init(value: $0) } + meansNumber = net.means.map { + NSNumber.init(value: $0) + } + dimsNum = [NSNumber.init(value: net.inputDim[0]), NSNumber.init(value: net.inputDim[3]), NSNumber.init(value: net.inputDim[1]), @@ -119,6 +122,29 @@ class ScaleKernel: CusomKernel { scaleTexture(input: texture!, complete: getTexture) } + @objc public func getTexture(inBuffer: MTLBuffer, getTexture: @escaping (MTLTexture) -> Void) { + guard let inQueue = queue, let inDevice = device else { + fatalError( " queue or devcie nil " ) + } + + guard let buffer = inQueue.makeCommandBuffer() else { + fatalError( " make buffer error" ) + } + + let bufferToTextureKernel = BufferToTextureKernel.init(device: inDevice, outputDim: Shape.init(inWidth: net.inputDim[2], inHeight: net.inputDim[1], inChannel: net.inputDim[3])) + do { + try bufferToTextureKernel.compute(inputBuffer: inBuffer, commandBuffer: buffer) + } catch { + fatalError(" bufferToTextureKernel error ") + } + + buffer.addCompletedHandler { (buffer) in + getTexture(bufferToTextureKernel.outputTexture) + } + + buffer.commit() + } + public func scaleTexture(input: MTLTexture , complete: @escaping (MTLTexture) -> Void) { guard let inQueue = queue, let inDevice = device else { @@ -129,7 +155,7 @@ class ScaleKernel: CusomKernel { fatalError( " make buffer error" ) } - let scaleKernel = ScaleKernel.init(device: inDevice, shape: CusomKernel.Shape.init(inWidth: net.inputDim[2], inHeight: net.inputDim[1], inChannel: 3)) + let scaleKernel = ScaleKernel.init(device: inDevice, shape: Shape.init(inWidth: net.inputDim[2], inHeight: net.inputDim[1], inChannel: 3)) do { try scaleKernel.compute(inputTexuture: input, commandBuffer: buffer) diff --git a/metal/paddle-mobile/paddle-mobile/framework/Executor.swift b/metal/paddle-mobile/paddle-mobile/framework/Executor.swift index 5c33fb9ea7..53c2690c5e 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Executor.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Executor.swift @@ -25,8 +25,7 @@ public class GPUResultHolder { public let capacity: Int public var resultPointer: UnsafeMutablePointer? public var intermediateResults: [String : [Variant]]? - public let elapsedTime: Double - public init(inDim: [Int], inPointer: UnsafeMutablePointer?, inCapacity: Int, inElapsedTime: Double, inIntermediateResults: [String : [Variant]]? = nil) { + public init(inDim: [Int], inPointer: UnsafeMutablePointer?, inCapacity: Int, inIntermediateResults: [String : [Variant]]? = nil) { dim = inDim capacity = inCapacity @@ -35,7 +34,6 @@ public class GPUResultHolder { resultPointer?.initialize(from: inInPointer, count: inCapacity) } - elapsedTime = inElapsedTime intermediateResults = inIntermediateResults } @@ -124,7 +122,6 @@ public class Executor { resInput = input } - let beforeDate = Date.init() let inputTexture = InputTexture.init(inMTLTexture: resInput, inExpectDim: dim) program.scope.setInput(input: inputTexture) //(ops.count - except) @@ -150,28 +147,28 @@ public class Executor { //将输入写进文件 /* + let inputArr = resInput.toTensor(dim: (n: dim[0], c: dim[3], h: dim[1], w: dim[2])) print(dim) - writeToLibrary(fileName: "test_image_yolo", array: inputArr) + writeToLibrary(fileName: "test_image_mingren", array: inputArr) print(" write done ") + return */ - /* 输出 op 计算结果 for op in SSelf.ops { op.delogOutput() } */ - let afterDate = Date.init() var resultHolder: GPUResultHolder if except > 0 { - resultHolder = GPUResultHolder.init(inDim: [], inPointer: nil, inCapacity: 0, inElapsedTime: afterDate.timeIntervalSince(beforeDate), inIntermediateResults: outputTextures) + resultHolder = GPUResultHolder.init(inDim: [], inPointer: nil, inCapacity: 0, inIntermediateResults: outputTextures) } else { let outputVar: Variant = SSelf.program.scope.output()! let output: FetchHolder = outputVar as! FetchHolder - resultHolder = GPUResultHolder.init(inDim: output.dim.dims, inPointer: output.result, inCapacity: output.capacity, inElapsedTime: afterDate.timeIntervalSince(beforeDate)) + resultHolder = GPUResultHolder.init(inDim: output.dim.dims, inPointer: output.result, inCapacity: output.capacity) } completionHandle(resultHolder) diff --git a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift index 462aff4b58..06c54290e6 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift @@ -96,8 +96,6 @@ public class Texture: Tensorial { return metalTexture.realNHWC(dim: (n: padToFourDim[0], h: padToFourDim[1], w: padToFourDim[2], c: padToFourDim[3])) } - - func initTexture(device: MTLDevice, inTranspose: [Int] = [0, 1, 2, 3], computePrecision: ComputePrecision = .Float16) { transpose = inTranspose for i in 0..<(4 - tensorDim.cout()) { diff --git a/src/framework/executor.cpp b/src/framework/executor.cpp index 751d04fc65..26e6877267 100644 --- a/src/framework/executor.cpp +++ b/src/framework/executor.cpp @@ -507,7 +507,7 @@ void Executor::Predict_To(int end) { #ifdef PADDLE_MOBILE_CL template <> void Executor::InitNoPersistableMemory( - const LoDTensor &input_tensor) { + const Tensor &input_tensor) { DLOG << "CL InitNoPersistableMemory "; for (const auto &block : program_desc_->Blocks()) { for (const auto &var_desc : block->Vars()) { -- GitLab