From 7db81656d8e19a64e6dddf8751a1b40e71a31760 Mon Sep 17 00:00:00 2001 From: liuruilong Date: Sat, 18 Aug 2018 14:20:53 +0800 Subject: [PATCH] operator init texture itself --- .../paddle-mobile-demo/ViewController.swift | 2 +- .../ViewController.swift | 1 - .../paddle-mobile/Common/Types.swift | 135 +++++++++++++++++- .../paddle-mobile/paddle-mobile/Loader.swift | 4 +- .../Kernels/ConvAddBatchNormReluKernel.swift | 4 +- .../Operators/Kernels/ConvAddKernel.swift | 3 + .../Operators/Kernels/Kernels.metal | 76 ++++++++++ .../paddle-mobile/Program/TensorDesc.swift | 10 +- .../paddle-mobile/framework/Tensor.swift | 10 +- .../paddle-mobile/framework/Texture.swift | 81 +++++------ 10 files changed, 264 insertions(+), 62 deletions(-) diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift index 241b50e2ef..f6d8d6c7c5 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift @@ -28,7 +28,7 @@ class ViewController: UIViewController { var selectImage: UIImage? var program: Program? var executor: Executor? - var modelType: SupportModel = .mobilenet + var modelType: SupportModel = SupportModel.supportedModels()[0] var toPredictTexture: MTLTexture? var modelHelper: Net { return modelHelperMap[modelType] ?! " has no this type " diff --git a/metal/paddle-mobile-unit-test/paddle-mobile-unit-test/ViewController.swift b/metal/paddle-mobile-unit-test/paddle-mobile-unit-test/ViewController.swift index d57b610e4d..c0be8996f1 100644 --- a/metal/paddle-mobile-unit-test/paddle-mobile-unit-test/ViewController.swift +++ b/metal/paddle-mobile-unit-test/paddle-mobile-unit-test/ViewController.swift @@ -10,7 +10,6 @@ import UIKit import paddle_mobile class ViewController: UIViewController { - override func viewDidLoad() { super.viewDidLoad() print(" done ") diff --git a/metal/paddle-mobile/paddle-mobile/Common/Types.swift b/metal/paddle-mobile/paddle-mobile/Common/Types.swift index 98353617f5..46f388d998 100644 --- a/metal/paddle-mobile/paddle-mobile/Common/Types.swift +++ b/metal/paddle-mobile/paddle-mobile/Common/Types.swift @@ -81,11 +81,140 @@ extension Float32: PrecisionType { } } -public enum DataLayout { - case NCHW - case NHWC +// N - 0 C - 1 H - 2 W - 3 +struct DataLayout { + + static func NCHW(dim: Dim = Dim.init(inDim: [0, 0, 0, 0])) -> DataLayout { + return DataLayout.init([(.N, dim[0]), (.C, dim[1]), (.H, dim[2]), (.W, dim[3])]) + } + + static func NHWC(dim: Dim = Dim.init(inDim: [0, 0, 0, 0])) -> DataLayout { + return DataLayout.init([(.N, dim[0]), (.H, dim[1]), (.W, dim[2]), (.C, dim[3])]) + } + + func count() -> Int { + return layoutWithDim.count + } + + var N: Int? { + get { + for layoutDim in layoutWithDim { + if layoutDim.0 == .N { + return layoutDim.1 + } + } + return nil + } + set { + var newN = (Layout.N, newValue) + if let index = layoutWithDim.index(where: { (layout: Layout, dim: Int) -> Bool in + return layout == .N + }) { + fatalError() + } + } + } + var C: Int? { + get { + for layoutDim in layoutWithDim { + if layoutDim.0 == .C { + return layoutDim.1 + } + } + return nil + } + set { + var newN = (Layout.C, newValue) + if let index = layoutWithDim.index(where: { (layout: Layout, dim: Int) -> Bool in + return layout == .N + }) { + fatalError() + } + } + } + var H: Int? { + get { + for layoutDim in layoutWithDim { + if layoutDim.0 == .H { + return layoutDim.1 + } + } + return nil + } + set { + var newN = (Layout.H, newValue) + if let index = layoutWithDim.index(where: { (layout: Layout, dim: Int) -> Bool in + return layout == .H + }) { + fatalError() + } + } + } + var W: Int? { + get { + for layoutDim in layoutWithDim { + if layoutDim.0 == .W { + return layoutDim.1 + } + } + return nil + } + set { + var newN = (Layout.W, newValue) + if let index = layoutWithDim.index(where: { (layout: Layout, dim: Int) -> Bool in + return layout == .W + }) { + fatalError() + } + } + } + + + init(_ inLayout: [(Layout, Int)]) { + layoutWithDim = inLayout + } + + func layout() -> [Layout] { + return layoutWithDim.map({ (layout: Layout, dim: Int) -> Layout in + return layout + }) + } + + var layoutWithDim: [(Layout, Int)] = [(.N, 0), (.C, 0), (.H, 0), (.W, 0)] + + func convertTo(inLayout: [Layout]) { + + } + + enum Layout: Int{ + case N = 0 + case C = 1 + case H = 2 + case W = 3 + static func defaultLayout() -> [Layout] { + return [N, C, H, W] + } + } } +extension DataLayout: Equatable { + public static func == (lhs: DataLayout, rhs: DataLayout) -> Bool { + if lhs.layoutWithDim.count == rhs.layoutWithDim.count { + var result = true + for i in 0.. { } catch let error { throw error } - tensor.convert(to: .NHWC) + tensor.convert(to: DataLayout.NHWC()) // tensor.initBuffer(device: device) scope[varDesc.name] = tensor } else { - let dim = Dim.init(inDim: tensorDesc.NHWCDim) + let dim = Dim.init(inDim: tensorDesc.dims) scope[varDesc.name] = Texture

.init(device: device, inDim: dim) } } else { diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift index 0ffe90272f..c7178e2f6f 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift @@ -50,6 +50,8 @@ class ConvAddBatchNormReluKernel: Kernel, Computable, Testable required init(device: MTLDevice, param: ConvAddBatchNormReluParam

) { + param.output.initTexture(device: device, transpose: [0, 2, 3, 1]) + if param.filter.width == 1 && param.filter.height == 1 { super.init(device: device, inFunctionName: "conv_add_batch_norm_relu_1x1") } else if param.filter.channel == 1 { @@ -60,12 +62,12 @@ class ConvAddBatchNormReluKernel: Kernel, Computable, Testable param.filter.initBuffer(device: device, precision: Tensor.BufferPrecision.Float32) param.y.initBuffer(device: device, precision: Tensor.BufferPrecision.Float32) - param.variance.initBuffer(device: device) param.mean.initBuffer(device: device) param.scale.initBuffer(device: device) param.bias.initBuffer(device: device) + let offsetX = param.filter.width/2 - Int(param.paddings[0]) let offsetY = param.filter.height/2 - Int(param.paddings[1]) diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift index 81f3aacba8..1f49fd14c7 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift @@ -13,6 +13,7 @@ limitations under the License. */ import Foundation +import MetalPerformanceShaders class ConvAddKernel: Kernel, Computable { var metalParam: MetalConvParam! @@ -32,6 +33,8 @@ class ConvAddKernel: Kernel, Computable { } func compute(commandBuffer: MTLCommandBuffer, param: ConvAddParam

) throws { + + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { throw PaddleMobileError.predictError(message: " encode is nil") } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal index 92ee118452..668ed443be 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal @@ -250,3 +250,79 @@ kernel void softmax_half(texture2d_array inTexture [[texture rr = exp(rr - maxv) / sum; outTexture.write(rr, gid.xy, gid.z); } + +kernel void prior_box(texture2d_array inTexture [[texture(0)]], + texture2d_array outTexture [[texture(1)]], + uint3 gid [[thread_position_in_grid]]) { + + int max_sizes_size; + float max_sizes[2]; + + bool clip; + + float img_width; + float img_height; + + float step_width; + float step_height; + float offset; + + float aspect_ratios[2]; + int aspect_ratios_size; + + float center_x = (gid.x + offset) * step_width; + float center_y = (gid.y + offset) * step_width; + + float box_width, box_height; + + int min_sizes_size; + float min_sizes[2]; + + float min_size; + float max_size; + + if (gid.z < aspect_ratios_size) { + float ar = aspect_ratios[gid.z]; + box_width = min_size * sqrt(ar) / 2; + box_height = min_size / sqrt(ar) / 2; + float4 box; + box.x = (center_x - box_width) / img_width; + box.y = (center_y - box_height) / img_height; + box.z = (center_x + box_width) / img_width; + box.w = (center_y + box_height) / img_height; + + float4 res; + if (clip) { + res = min(max(box, 0.0), 1.0); + } else { + res = box; + } + + outTexture.write(res, gid.xy, gid.z); + } else if (gid.z >= aspect_ratios_size) { + int max_index = gid.z - aspect_ratios_size; + if (max_sizes_size > 0 && min_sizes_size > 0) { + box_width = box_height = sqrt(min_size * max_size) / 2; + float4 max_box; + max_box.x = (center_x - box_width) / img_width; + max_box.y = (center_y - box_height) / img_height; + max_box.z = (center_x + box_width) / img_width; + max_box.w = (center_y + box_height) / img_height; + + float4 res; + if (clip) { + res = min(max(max_box, 0.0), 1.0); + } else { + res = max_box; + } + + outTexture.write(max_box, gid.xy, gid.z); + } + } +} + + + + + + diff --git a/metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift b/metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift index e564821ab6..be37b85e1f 100644 --- a/metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift +++ b/metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift @@ -17,15 +17,15 @@ import Foundation struct TensorDesc { let dims: [Int] let dataType: VarTypeType - let dataLayout: DataLayout = .NCHW + let dataLayout: DataLayout = DataLayout.NHWC() var NCHWDim: [Int] { get { if dims.count != 4 { return dims } - if dataLayout == .NCHW { + if dataLayout == DataLayout.NCHW() { return dims - } else if dataLayout == .NHWC{ + } else if dataLayout == DataLayout.NHWC() { var resultDims = dims resultDims.swapAt(1, 3) return resultDims @@ -40,9 +40,9 @@ struct TensorDesc { if dims.count != 4 { return dims } - if dataLayout == .NHWC { + if dataLayout == DataLayout.NHWC() { return dims - } else if dataLayout == .NCHW{ + } else if dataLayout == DataLayout.NCHW() { var resultDims = dims resultDims.swapAt(1, 3) return resultDims diff --git a/metal/paddle-mobile/paddle-mobile/framework/Tensor.swift b/metal/paddle-mobile/paddle-mobile/framework/Tensor.swift index 7ffcd97f44..d2c8228ed9 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Tensor.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Tensor.swift @@ -61,7 +61,7 @@ class Tensor: Tensorial { } } - required init(inDim: Dim, inLayout: DataLayout = .NCHW) { + required init(inDim: Dim, inLayout: DataLayout = DataLayout.NCHW()) { dim = inDim let size = inDim.numel() * MemoryLayout

.size let pointer = UnsafeMutablePointer

.allocate(capacity: size) @@ -78,13 +78,13 @@ class Tensor: Tensorial { return } - guard layout == .NCHW && to == .NHWC else { + guard layout == DataLayout.NCHW() && to == DataLayout.NHWC() else { // other not support return } let newPointer = UnsafeMutablePointer

.allocate(capacity: data.size) - if layout == .NCHW { + if layout == DataLayout.NCHW() { NCHW2NHWC(newPtr: newPointer) } @@ -106,7 +106,6 @@ class Tensor: Tensorial { fatalError(" not support yet ") } - let precisionSize: Int switch precision { case .Float32: @@ -116,7 +115,7 @@ class Tensor: Tensorial { } if dim.cout() == 4 { - if layout == .NHWC { + if layout == DataLayout.NHWC() { let C = dim[3] let cSlices = (C + 3) / 4 let paddedC = cSlices * 4 @@ -232,7 +231,6 @@ class Tensor: Tensorial { } } - extension Tensor { var debugDescription: String { diff --git a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift index d217d74abf..c26994116d 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift @@ -40,59 +40,34 @@ extension InputTexture { public class Texture: Tensorial { var dim: Dim - let textureDesc: MTLTextureDescriptor - var metalTexture: MTLTexture + private(set) var originDim: Dim + private var textureDesc: MTLTextureDescriptor! + var metalTexture: MTLTexture! + var transpose: [Int] = [0, 1, 2, 3] - init(device: MTLDevice, inDim: Dim, inLayout: DataLayout = .NHWC) { - dim = inDim - layout = inLayout + func initTexture(device: MTLDevice, transpose: [Int]) { + let newDim = transpose.map { originDim[$0] } + + let newLayout = transpose.map {layout.layoutWithDim[$0] } + + layout = DataLayout.init(newLayout) + dim = Dim.init(inDim: newDim) + let tmpTextureDes = MTLTextureDescriptor.init() - if inDim.cout() == 1 { - tmpTextureDes.width = inDim[0] - tmpTextureDes.textureType = .type1D - } else if inDim.cout() == 4 { - tmpTextureDes.height = inDim[1] - tmpTextureDes.width = inDim[2] - tmpTextureDes.depth = 1 - tmpTextureDes.arrayLength = (inDim[3] * inDim[0] + 3)/4 - tmpTextureDes.textureType = .type2DArray - } else if inDim.cout() == 2 { -// tmpTextureDes.height = 1 -// tmpTextureDes.width = 1 -// tmpTextureDes.depth = 1 -// tmpTextureDes.arrayLength = (inDim[0] * inDim[1] + 3)/4 - - tmpTextureDes.width = inDim[0] - tmpTextureDes.height = inDim[1] - tmpTextureDes.depth = 1 - tmpTextureDes.arrayLength = 1 - tmpTextureDes.textureType = .type2DArray - } else { - /* - var name: box_coder_0.tmp_0 - in var tensor desc dims size: 3 - var tensor desc dim 0 value: -1 - var tensor desc dim 1 value: 1917 - var tensor desc dim 2 value: 4 - */ - - tmpTextureDes.height = inDim[1] - tmpTextureDes.width = inDim[2] - tmpTextureDes.depth = 1 - tmpTextureDes.arrayLength = 1 - tmpTextureDes.textureType = .type2DArray - } + + tmpTextureDes.width = layout.W ?? 1 + tmpTextureDes.height = layout.H ?? 1 + tmpTextureDes.depth = 1 + tmpTextureDes.arrayLength = ((layout.N ?? 1) * (layout.C ?? 1) + 3) / 4 + tmpTextureDes.textureType = .type2DArray if MemoryLayout

.size == 1 { tmpTextureDes.pixelFormat = .rgba8Unorm } else if MemoryLayout

.size == 2 { tmpTextureDes.pixelFormat = .rgba16Float } else if MemoryLayout

.size == 4 { -// tmpTextureDes.pixelFormat = .r32Float tmpTextureDes.pixelFormat = .rgba32Float - } -// tmpTextureDes.pixelFormat = .rgba16Float tmpTextureDes.usage = [.shaderRead, .shaderWrite] tmpTextureDes.storageMode = .shared @@ -100,6 +75,26 @@ public class Texture: Tensorial { metalTexture = device.makeTexture(descriptor: tmpTextureDes) ?! " texture nil " } + init(device: MTLDevice, inDim: Dim) { + var fourDim: Dim + if inDim.cout() == 4 { + fourDim = inDim + } else if inDim.cout() < 4 { + var fourDimNum: [Int] = [] + for _ in 0..<(4 - inDim.cout()) { + fourDimNum.append(1) + } + fourDimNum.append(contentsOf: inDim.dims) + fourDim = Dim.init(inDim: fourDimNum) + } else { + fatalError(" not support ") + } + + dim = fourDim + originDim = fourDim + layout = DataLayout.init([(.N, fourDim[0]), (.C, fourDim[1]), (.H, fourDim[2]), (.W, fourDim[3])]) + } + // required public init(inDim: Dim, inLayout: DataLayout = .NHWC, inTexture: MTLTexture) { // dim = inDim // layout = inLayout -- GitLab