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 6f3db873591d1d4ca05fc1ea563d01f44c779d7e..6d057200279c3d2472c764587e5a93dd0352526e 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj +++ b/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj @@ -16,6 +16,7 @@ FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */; }; 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 */; }; @@ -56,6 +57,7 @@ FC039B8D20E11C560081E9F8 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; sourceTree = ""; }; FC918190211DBC3500B6F354 /* paddle-mobile.png */ = {isa = PBXFileReference; lastKnownFileType = image.png; path = "paddle-mobile.png"; sourceTree = ""; }; FC918192211DC70500B6F354 /* iphone.JPG */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = iphone.JPG; sourceTree = ""; }; + FCA3A16021313E1F00084FE5 /* hand.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = hand.jpg; sourceTree = ""; }; FCBCCC502122EEDC00D94F7E /* ssd_hand_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_params; sourceTree = ""; }; FCBCCC512122EEDC00D94F7E /* ssd_hand_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_model; sourceTree = ""; }; FCBCCC542122EF5400D94F7E /* MetalHelper.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MetalHelper.swift; sourceTree = ""; }; @@ -137,6 +139,7 @@ FC0E2C1D20EDC030009C1FAC /* images */ = { isa = PBXGroup; children = ( + FCA3A16021313E1F00084FE5 /* hand.jpg */, FC918192211DC70500B6F354 /* iphone.JPG */, FC918190211DBC3500B6F354 /* paddle-mobile.png */, FCDFD41A211D91C7005AB38B /* synset.txt */, @@ -245,6 +248,7 @@ FCDFD41B211D91C7005AB38B /* synset.txt in Resources */, FCD04E6420F3146B0007374F /* model in Resources */, FC039B8720E11C550081E9F8 /* Main.storyboard in Resources */, + FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */, FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */, ); runOnlyForDeploymentPostprocessing = 0; diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/Base.lproj/Main.storyboard b/metal/paddle-mobile-demo/paddle-mobile-demo/Base.lproj/Main.storyboard index a5efadeb97ccc41449dc32a2c1dfcdfcf9fceac5..38094021c889611792b0041abf55b81f1dc04966 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/Base.lproj/Main.storyboard +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/Base.lproj/Main.storyboard @@ -19,10 +19,10 @@ - + - - + @@ -203,6 +203,7 @@ + diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift index 86dda2c5902da1473ded2af1c443a07335ffafe1..6ff4bbe8e58e2b245e581b7519bc6fc86cfee8cf 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/ModelHelper.swift @@ -30,6 +30,7 @@ protocol Net { var preprocessKernel: CusomKernel { get } func getTexture(image: CGImage, getTexture: @escaping (MTLTexture) -> Void) func resultStr(res: [Float]) -> String + func fetchResult(paddleMobileRes: ResultHolder) -> [Float32] } extension Net { @@ -39,10 +40,13 @@ extension Net { getTexture(resTexture) } } + + func fetchResult(paddleMobileRes: ResultHolder) -> [Float32] { + return paddleMobileRes.resultArr + } } struct MobileNet: Net{ - class MobilenetPreProccess: CusomKernel { init(device: MTLDevice) { let s = CusomKernel.Shape.init(inWidth: 224, inHeight: 224, inChannel: 3) @@ -100,7 +104,8 @@ struct MobileNet_ssd_hand: Net{ } func resultStr(res: [Float]) -> String { - fatalError() + return "哈哈哈, 还没好" +// fatalError() } func bboxArea(box: [Float32], normalized: Bool) -> Float32 { @@ -117,7 +122,6 @@ struct MobileNet_ssd_hand: Net{ } } - func jaccardOverLap(box1: [Float32], box2: [Float32], normalized: Bool) -> Float32 { if box2[0] > box1[2] || box2[2] < box1[0] || box2[1] > box1[3] || box2[3] < box1[1] { @@ -136,9 +140,11 @@ struct MobileNet_ssd_hand: Net{ } } - func fetchResult(paddleMobileRes: [String : Texture]) -> [Float32]{ - let bbox = paddleMobileRes["box_coder_0.tmp_0"] ?! " no bbox " - let scores = paddleMobileRes["transpose_12.tmp_0"] ?! " no scores " + func fetchResult(paddleMobileRes: ResultHolder) -> [Float32]{ + let scores = paddleMobileRes.intermediateResults![0] as! Texture + let bbox = paddleMobileRes.intermediateResults![1] as! Texture +// let bbox = paddleMobileRes["box_coder_0.tmp_0"] ?! " no bbox " +// let scores = paddleMobileRes["transpose_12.tmp_0"] ?! " no scores " let score_thredshold: Float32 = 0.01 let nms_top_k = 400 let keep_top_k = 200 @@ -156,20 +162,29 @@ struct MobileNet_ssd_hand: Net{ var scoreFormatArr: [Float32] = [] var outputArr: [Float32] = [] - let numOfOneC = (scores.originDim[2] + 3) / 4 // 480 - let cNumOfOneClass = numOfOneC * 4 // 1920 + let numOfOneC = (scores.tensorDim[2] + 3) / 4 // 480 - let boxSize = bbox.originDim[2] // 4 - let classNum = scores.originDim[1] // 7 + let cNumOfOneClass = scores.tensorDim[2] // 1917 + + let cPaddedNumOfOneClass = numOfOneC * 4 // 1920 + + let boxSize = bbox.tensorDim[2] // 4 + let classNum = scores.tensorDim[1] // 7 let classNumOneTexture = classNum * 4 // 28 for c in 0..(scoreFormatArr[(i * cNumOfOneClass)..<((i + 1) * cNumOfOneClass)]) var scoreThresholdArr: [(Float32, Int)] = [] - for i in 0.. score_thredshold { - scoreThresholdArr.append((sliceScore[i], i)) + for j in 0.. score_thredshold { + scoreThresholdArr.append((sliceScore[j], j)) } } @@ -204,7 +219,7 @@ struct MobileNet_ssd_hand: Net{ if keep { let keptIdx = selectedIndex[j].0 let box1 = Array(bboxArr[(idx * boxSize)..<(idx * boxSize + 4)]) - let box2 = Array(bboxArr[(idx * boxSize)..<(keptIdx * boxSize + 4)]) + let box2 = Array(bboxArr[(keptIdx * boxSize)..<(keptIdx * boxSize + 4)]) let overlap = jaccardOverLap(box1: box1, box2: box2, normalized: true) keep = (overlap <= nms_threshold) @@ -259,7 +274,8 @@ struct MobileNet_ssd_hand: Net{ outputArr.append(contentsOf: subBox) } } - + print(" fuck success !") + print(outputArr) return outputArr } diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift index fcf6a6cd5fa2c6552f173cf2e83edbe1a20c1fb6..37151b3ba7819177ed99d5653dfb218752856d7a 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift @@ -75,7 +75,7 @@ class ViewController: UIViewController { } do { - let max = 10 + let max = 1 var startDate = Date.init() for i in 0..(option: T?, excuteOrError: @autoclosure () -> String) -> T{ - if let inOpt = option { - return inOpt - }else{ - print(excuteOrError()) - fatalError(excuteOrError()) - } + if let inOpt = option { + return inOpt + }else{ + print(excuteOrError()) + fatalError(excuteOrError()) + } } //Lense struct Lense { - let from: (A) -> B - let to: (B, A) -> A + let from: (A) -> B + let to: (B, A) -> A } precedencegroup CombineLense{ - associativity: left - higherThan: AssignmentPrecedence + associativity: left + higherThan: AssignmentPrecedence } infix operator >>>: CombineLense func >>>(left: Lense, right: Lense) -> Lense { - return Lense.init(from: { (a) -> C in - left.from(right.from(a)) - }, to: { (c, a) -> A in - right.to( left.to(c, right.from(a)),a) - }) + return Lense.init(from: { (a) -> C in + left.from(right.from(a)) + }, to: { (c, a) -> A in + right.to( left.to(c, right.from(a)),a) + }) } protocol CIntIndex { - associatedtype T; - subscript(index: CInt) -> T { get set}; + associatedtype T; + subscript(index: CInt) -> T { get set}; } extension Array: CIntIndex{ - typealias T = Element - subscript(index: CInt) -> T { - get{ - guard Int64(Int.max) >= Int64(index) else{ - fatalError("cint index out of Int range") - } - return self[Int(index)] - } - set{ - guard Int64(Int.max) >= Int64(index) else{ - fatalError("cint index out of Int range") - } - self[Int(index)] = newValue - } - + typealias T = Element + subscript(index: CInt) -> T { + get{ + guard Int64(Int.max) >= Int64(index) else{ + fatalError("cint index out of Int range") + } + return self[Int(index)] + } + set{ + guard Int64(Int.max) >= Int64(index) else{ + fatalError("cint index out of Int range") + } + self[Int(index)] = newValue } + + } } extension Array where Element: AnyObject{ - mutating func remove(element: Element) { - if let index = index(where: { (node) -> Bool in - return unsafeBitCast(element, to: Int.self) == unsafeBitCast(node, to: Int.self) - }) { - remove(at: index) - } + mutating func remove(element: Element) { + if let index = index(where: { (node) -> Bool in + return unsafeBitCast(element, to: Int.self) == unsafeBitCast(node, to: Int.self) + }) { + remove(at: index) } - + } + } //MARK: Array extension extension Array where Element: Comparable{ - - /// 返回数组前 r 个元素, 并将元素处于原数组的位置作为元组的第一个元素返回 - /// - /// - Parameter r: 前 r 个元素 - /// - Returns: [(原有位置, 排好位置的元素)] - public func top(r: Int) -> [(Int, Element)] { - precondition(r <= self.count) - return Array<(Int, Element)>(zip(0.. $1.1 }.prefix(through: r - 1)) + + /// 返回数组前 r 个元素, 并将元素处于原数组的位置作为元组的第一个元素返回 + /// + /// - Parameter r: 前 r 个元素 + /// - Returns: [(原有位置, 排好位置的元素)] + public func top(r: Int) -> [(Int, Element)] { + precondition(r <= self.count) + return Array<(Int, Element)>(zip(0.. $1.1 }.prefix(through: r - 1)) + } +} + +extension Array { + func strideArray(inCount: Int = 20) -> Array { + if count < inCount { + return self + } else { + let stride = count / inCount + var newArray: [Element] = [] + for i in 0.. UnsafePointer? { - return (self as NSString).utf8String - } + func cStr() -> UnsafePointer? { + return (self as NSString).utf8String + } } func address(o: T) -> String { - return String.init(format: "%018p", unsafeBitCast(o, to: Int.self)) + return String.init(format: "%018p", unsafeBitCast(o, to: Int.self)) } diff --git a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift index 71b50f0553cbef8d1f0815f95c4ce9e4af150021..58295cfb5449718da0204cd2498287658c6f77cb 100644 --- a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift +++ b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift @@ -18,365 +18,372 @@ fileprivate var defaultMetalLibrary: MTLLibrary? fileprivate var paddleMobileMetalLibrary: MTLLibrary? extension MTLDevice { - func defaultLibrary() -> MTLLibrary { - if defaultMetalLibrary == nil { - defaultMetalLibrary = makeDefaultLibrary() - } - if let inDefaultLib = defaultMetalLibrary { - return inDefaultLib - } else { - fatalError(" default metal libary is nil") - } + func defaultLibrary() -> MTLLibrary { + if defaultMetalLibrary == nil { + defaultMetalLibrary = makeDefaultLibrary() } - - func paddleMobileLibrary() -> MTLLibrary { - if paddleMobileMetalLibrary == nil { - guard let path = Bundle.init(for: Kernel.self).path(forResource: "default", ofType: "metallib") else { - fatalError("Counld't find paddle mobile library") - } - do { - paddleMobileMetalLibrary = try makeLibrary(filepath: path) - } catch _ { - fatalError("Counld't load paddle mobile library") - } - } - - if let inPaddleMobileLib = paddleMobileMetalLibrary { - return inPaddleMobileLib - } else { - fatalError("PaddleMobile metal libary is nil") - } + if let inDefaultLib = defaultMetalLibrary { + return inDefaultLib + } else { + fatalError(" default metal libary is nil") } - - func pipeLine(funcName: String, inPaddleMobileLib: Bool = true) -> MTLComputePipelineState { - let useLib = inPaddleMobileLib ? paddleMobileLibrary() : defaultLibrary() - guard let function = useLib.makeFunction(name: funcName) else { - fatalError(" function " + funcName + " not found") - } - do { - let pipLine = try makeComputePipelineState(function: function) - return pipLine - } catch _ { - fatalError("make pip line error occured") - } - + } + + func paddleMobileLibrary() -> MTLLibrary { + if paddleMobileMetalLibrary == nil { + guard let path = Bundle.init(for: Kernel.self).path(forResource: "default", ofType: "metallib") else { + fatalError("Counld't find paddle mobile library") + } + do { + paddleMobileMetalLibrary = try makeLibrary(filepath: path) + } catch _ { + fatalError("Counld't load paddle mobile library") + } } - func makeBuffer

(value: [P]) -> MTLBuffer { - let buffer = makeBuffer(length: value.count * MemoryLayout

.size, options: MTLResourceOptions.storageModeShared) - let contents = buffer?.contents().bindMemory(to: P.self, capacity: value.count * MemoryLayout

.size) - for i in 0.. MTLComputePipelineState { + let useLib = inPaddleMobileLib ? paddleMobileLibrary() : defaultLibrary() + guard let function = useLib.makeFunction(name: funcName) else { + fatalError(" function " + funcName + " not found") + } + do { + let pipLine = try makeComputePipelineState(function: function) + return pipLine + } catch _ { + fatalError("make pip line error occured") } - func texture2tensor

(texture: MTLTexture, dim: [Int], transpose: [Int] = [0, 1, 2, 3]) -> [P] { - var tdim: [Int] = [1, 1, 1, 1] - for i in 0...size - let bpI = ndim[1] * bpR - let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: ndim[2], height: ndim[1], depth: 1)) - for i in 0.. = UnsafeMutablePointer

.allocate(capacity: ndim[1] * ndim[2] * 4 * MemoryLayout

.size) - texture.getBytes(pointer, bytesPerRow: bpR, bytesPerImage: bpI, from: region, mipmapLevel: 0, slice: i) - - for h in 0..(value: [P]) -> MTLBuffer { + let buffer = makeBuffer(length: value.count * MemoryLayout

.size, options: MTLResourceOptions.storageModeShared) + let contents = buffer?.contents().bindMemory(to: P.self, capacity: value.count * MemoryLayout

.size) + for i in 0..(texture: MTLTexture, dim: [Int], transpose: [Int] = [0, 1, 2, 3]) -> [P] { + var tdim: [Int] = [1, 1, 1, 1] + for i in 0..(value: [P], dim: [Int], transpose: [Int] = [0, 1, 2, 3]) -> MTLTexture { - if value.count > 0 { - assert(value.count == dim.reduce(1) { $0 * $1 }) - } - - var tdim: [Int] = [1, 1, 1, 1] - for i in 0.. 0 { - var rcount: Int = (ndim[0] * ndim[3] + 3) / 4 - rcount = rcount * 4 * ndim[1] * ndim[2] - var nvalue: [P] = .init(repeating: Float32(0.0) as! P, count: rcount) - - for i0 in 0...size + let bpI = ndim[1] * bpR + let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: ndim[2], height: ndim[1], depth: 1)) + for i in 0.. = UnsafeMutablePointer

.allocate(capacity: ndim[1] * ndim[2] * 4 * MemoryLayout

.size) + texture.getBytes(pointer, bytesPerRow: bpR, bytesPerImage: bpI, from: region, mipmapLevel: 0, slice: i) + + for h in 0.. = UnsafeMutablePointer(mutating: nvalue) - let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: ndim[2], height: ndim[1], depth: 1)) - let bpR = ndim[2] * 4 * MemoryLayout

.size - let bpI = ndim[1] * bpR - for i in 0..(value: [P], textureWidth: Int, textureHeight: Int, arrayLength: Int) -> MTLTexture{ - - let textureDesc = MTLTextureDescriptor.init() - textureDesc.width = textureWidth - textureDesc.height = textureHeight - textureDesc.depth = 1 - textureDesc.usage = [.shaderRead, .shaderWrite] - textureDesc.pixelFormat = .rgba32Float - textureDesc.textureType = .type2DArray - textureDesc.storageMode = .shared - textureDesc.cpuCacheMode = .defaultCache - textureDesc.arrayLength = arrayLength - let texture = makeTexture(descriptor: textureDesc)! - - if value.count >= 4{ - let counts = arrayLength * 4 * textureWidth * textureHeight - let pointer: UnsafeMutablePointer

= UnsafeMutablePointer

.allocate(capacity: counts * MemoryLayout

.size) - for i in 0...size - let bytesPerImage = texture.height * bytesPerRow - let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: texture.width, height: texture.height, depth: texture.depth)) - for i in 0..(value: [P], dim: [Int], transpose: [Int] = [0, 1, 2, 3]) -> MTLTexture { + if value.count > 0 { + assert(value.count == dim.reduce(1) { $0 * $1 }) + } + + var tdim: [Int] = [1, 1, 1, 1] + for i in 0.. 0 { + var rcount: Int = (ndim[0] * ndim[3] + 3) / 4 + rcount = rcount * 4 * ndim[1] * ndim[2] + var nvalue: [P] = .init(repeating: Float32(0.0) as! P, count: rcount) + + for i0 in 0.. = UnsafeMutablePointer(mutating: nvalue) + let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: ndim[2], height: ndim[1], depth: 1)) + let bpR = ndim[2] * 4 * MemoryLayout

.size + let bpI = ndim[1] * bpR + for i in 0..(value: [P], textureWidth: Int, textureHeight: Int, arrayLength: Int) -> MTLTexture{ + + let textureDesc = MTLTextureDescriptor.init() + textureDesc.width = textureWidth + textureDesc.height = textureHeight + textureDesc.depth = 1 + textureDesc.usage = [.shaderRead, .shaderWrite] + textureDesc.pixelFormat = .rgba32Float + textureDesc.textureType = .type2DArray + textureDesc.storageMode = .shared + textureDesc.cpuCacheMode = .defaultCache + textureDesc.arrayLength = arrayLength + let texture = makeTexture(descriptor: textureDesc)! + + if value.count >= 4{ + let counts = arrayLength * 4 * textureWidth * textureHeight + let pointer: UnsafeMutablePointer

= UnsafeMutablePointer

.allocate(capacity: counts * MemoryLayout

.size) + for i in 0...size + let bytesPerImage = texture.height * bytesPerRow + let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: texture.width, height: texture.height, depth: texture.depth)) + for i in 0..(stridable: Bool = true) -> [(index: Int, value: P)] { - var arr: [P] = floatArray { (p: P) -> P in - return p; - } - var result: [(index: Int, value: P)] = [] - if arr.count > 100 && stridable { - for j in stride(from: 0, to: arr.count , by: arr.count / 100){ - result.append((j, arr[j])) - } - } else { - for j in 0..(stridable: Bool = true) -> [(index: Int, value: P)] { + var arr: [P] = floatArray { (p: P) -> P in + return p; } - - func floatArray(res: (P) -> T) -> [T] { - var fArr: [T] = [] - if textureType == .type2DArray { - for i in 0...size, alignment: MemoryLayout

.alignment) - let bytesPerRow = width * depth * 4 * MemoryLayout

.size - let bytesPerImage = width * height * depth * 4 * MemoryLayout

.size - let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: width, height: height, depth: depth)) - getBytes(bytes, bytesPerRow: bytesPerRow, bytesPerImage: bytesPerImage, from: region, mipmapLevel: 0, slice: i) - let p = bytes.assumingMemoryBound(to: P.self) - - for j in 0...size, alignment: MemoryLayout

.alignment) - let bytesPerRow = width * depth * 4 * MemoryLayout

.size - let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: width, height: height, depth: depth)) - getBytes(bytes, bytesPerRow: bytesPerRow, from: region, mipmapLevel: 0) - let p = bytes.assumingMemoryBound(to: P.self) - - for j in 0.. 100 && stridable { + for j in stride(from: 0, to: arr.count , by: arr.count / 100){ + result.append((j, arr[j])) + } + } else { + for j in 0..(res: (P) -> T) -> [T] { + var fArr: [T] = [] + if textureType == .type2DArray { + for i in 0...size, alignment: MemoryLayout

.alignment) + let bytesPerRow = width * depth * 4 * MemoryLayout

.size + let bytesPerImage = width * height * depth * 4 * MemoryLayout

.size + let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: width, height: height, depth: depth)) + getBytes(bytes, bytesPerRow: bytesPerRow, bytesPerImage: bytesPerImage, from: region, mipmapLevel: 0, slice: i) + let p = bytes.assumingMemoryBound(to: P.self) + + for j in 0...size, alignment: MemoryLayout

.alignment) + let bytesPerRow = width * depth * 4 * MemoryLayout

.size + let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: width, height: height, depth: depth)) + getBytes(bytes, bytesPerRow: bytesPerRow, from: region, mipmapLevel: 0) + let p = bytes.assumingMemoryBound(to: P.self) + + for j in 0..(header: String = "", stridable: Bool = true) -> T? { - print(header) - print("texture: \(self)") -// let res: [(index: Int, value: T)] = stridableFloatArray(stridable: stridable) -// print(res) + return fArr + } - if textureType == .type2DArray { - for i in 0...size, alignment: MemoryLayout.alignment) - let bytesPerRow = width * depth * 4 * MemoryLayout.size - let bytesPerImage = width * height * depth * 4 * MemoryLayout.size - let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: width, height: height, depth: depth)) - 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 && 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])" - } - } else { - for j in 0...size, alignment: MemoryLayout.alignment) - let bytesPerRow = width * depth * 4 * MemoryLayout.size - let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: width, height: height, depth: depth)) - getBytes(bytes, bytesPerRow: bytesPerRow, from: region, mipmapLevel: 0) - let p = bytes.assumingMemoryBound(to: T.self) - str += "2d count : \(width * width * 4) \n" - - if stridable { - for j in stride(from: 0, to: width * height * 4, by: width * height * 4 / 100){ - str += "index \(j): \(p[j]) " - } - } else { - for j in 0..(header: String = "", stridable: Bool = true) -> T? { + print(header) + print("texture: \(self)") + // let res: [(index: Int, value: T)] = stridableFloatArray(stridable: stridable) + // print(res) + + if textureType == .type2DArray { + for i in 0...size, alignment: MemoryLayout.alignment) + let bytesPerRow = width * depth * 4 * MemoryLayout.size + let bytesPerImage = width * height * depth * 4 * MemoryLayout.size + let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: width, height: height, depth: depth)) + 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 && width * height * depth * 4 > 20 { + for j in stride(from: 0, to: width * height * depth * 4 , by: width * height * depth * 4 / 20){ + str += " index \(j): \(p[j])" + } + } else { + for j in 0...size, alignment: MemoryLayout.alignment) + let bytesPerRow = width * depth * 4 * MemoryLayout.size + let region = MTLRegion.init(origin: MTLOrigin.init(x: 0, y: 0, z: 0), size: MTLSize.init(width: width, height: height, depth: depth)) + getBytes(bytes, bytesPerRow: bytesPerRow, from: region, mipmapLevel: 0) + let p = bytes.assumingMemoryBound(to: T.self) + str += "2d count : \(width * width * 4) \n" + + if stridable { + for j in stride(from: 0, to: width * height * 4, by: width * height * 4 / 20){ + str += "index \(j): \(p[j]) " + } + } else { + for j in 0..(header: String = "", stridable: Bool = true) -> T? { - print(header) - print("MTLBuffer: \(self) ") - var str = "" - if stridable && length/MemoryLayout.stride > 1000{ - for j in stride(from: 0, to: length, by: length/MemoryLayout.stride / 100){ - str += " \(contents().assumingMemoryBound(to: T.self)[j])" - } - } else { - for i in 0...size { - str += " \(contents().assumingMemoryBound(to: T.self)[i])" - } - } - print(str) - return nil + func logDesc(header: String = "", stridable: Bool = true) -> T? { + print(header) + print("MTLBuffer: \(self) ") + var str = "" + if stridable && length/MemoryLayout.stride > 1000{ + for j in stride(from: 0, to: length, by: length/MemoryLayout.stride / 100){ + str += " \(contents().assumingMemoryBound(to: T.self)[j])" + } + } else { + for i in 0...size { + str += " \(contents().assumingMemoryBound(to: T.self)[i])" + } } - - func makeTexture(textureWidth: Int, textureHeight: Int, arrayLength: Int) -> MTLTexture { - let textureDesc = MTLTextureDescriptor.init() - textureDesc.width = textureWidth - textureDesc.height = textureHeight - textureDesc.depth = 1 - textureDesc.usage = [.shaderRead, .shaderWrite] - textureDesc.pixelFormat = .rgba32Float - textureDesc.textureType = .type2DArray - textureDesc.storageMode = .shared - textureDesc.cpuCacheMode = .defaultCache - textureDesc.arrayLength = arrayLength - let texture = makeTexture(descriptor: textureDesc, offset: 0, bytesPerRow: textureWidth * 4 * 4)! - return texture + print(str) + return nil + } + + func makeTexture(textureWidth: Int, textureHeight: Int, arrayLength: Int) -> MTLTexture { + let textureDesc = MTLTextureDescriptor.init() + textureDesc.width = textureWidth + textureDesc.height = textureHeight + textureDesc.depth = 1 + textureDesc.usage = [.shaderRead, .shaderWrite] + textureDesc.pixelFormat = .rgba32Float + textureDesc.textureType = .type2DArray + textureDesc.storageMode = .shared + textureDesc.cpuCacheMode = .defaultCache + textureDesc.arrayLength = arrayLength + let texture = makeTexture(descriptor: textureDesc, offset: 0, bytesPerRow: textureWidth * 4 * 4)! + return texture + } + + func array() -> [T] { + var array: [T] = [] + let pointer = contents().bindMemory(to: T.self, capacity: length) + for i in 0..<(length / MemoryLayout.size) { + array.append(pointer[i]) } - - - + return array; + } + } diff --git a/metal/paddle-mobile/paddle-mobile/Common/Types.swift b/metal/paddle-mobile/paddle-mobile/Common/Types.swift index 46f388d998d829039bd89b32bdac09f480857999..cb39d44e2125cb616218c4648ec792526b608a5d 100644 --- a/metal/paddle-mobile/paddle-mobile/Common/Types.swift +++ b/metal/paddle-mobile/paddle-mobile/Common/Types.swift @@ -15,207 +15,210 @@ import Foundation public protocol SummableMultipliable: Equatable { - static func +(lhs: Self, rhs: Self) -> Self - static func *(lhs: Self, rhs: Self) -> Self - static func -(lhs: Self, rhs: Self) -> Self + static func +(lhs: Self, rhs: Self) -> Self + static func *(lhs: Self, rhs: Self) -> Self + static func -(lhs: Self, rhs: Self) -> Self } public protocol PrecisionType: SummableMultipliable{ - init(inFloat: Float32) - init(inFloat16: Float16) - init(_ inP: P) - static var bitSize: UInt { get } + init(inFloat: Float32) + init(inFloat16: Float16) + init(_ inP: P) + static var bitSize: UInt { get } } public typealias Float16 = Int16 extension Float16: PrecisionType { - public static func * (prefix: Float16, postfix: Float16) { - return prefix * postfix - } - - public init

(_ inP: P) where P : PrecisionType { - if P.bitSize == Float32.bitSize { - self = Float16(inFloat: inP as! Float32) - } else if P.bitSize == Float16.bitSize { - self = inP as! Float16 - } else { - fatalError() - } - } - - public static var bitSize: UInt { - return 16 - } - - public init(inFloat16: Float16) { - self = inFloat16 - } - public init(inFloat: Float32) { - self = Int16(inFloat) - } - - - + public static func * (prefix: Float16, postfix: Float16) { + return prefix * postfix + } + + public init

(_ inP: P) where P : PrecisionType { + if P.bitSize == Float32.bitSize { + self = Float16(inFloat: inP as! Float32) + } else if P.bitSize == Float16.bitSize { + self = inP as! Float16 + } else { + fatalError() + } + } + + public static var bitSize: UInt { + return 16 + } + + public init(inFloat16: Float16) { + self = inFloat16 + } + public init(inFloat: Float32) { + self = Int16(inFloat) + } + + + } extension Float32: PrecisionType { - public init

(_ inP: P) where P : PrecisionType { - if P.bitSize == Float32.bitSize { - self = inP as! Float32 - } else if P.bitSize == Float16.bitSize { - self = Float32.init(inP as! Float16) - } else { - fatalError() - } - } - - public init(inFloat: Float32) { - self = inFloat - } - - public init(inFloat16: Float16) { - self = Float32.init(inFloat16) - } - - public static var bitSize: UInt { - return 32 - } + public init

(_ inP: P) where P : PrecisionType { + if P.bitSize == Float32.bitSize { + self = inP as! Float32 + } else if P.bitSize == Float16.bitSize { + self = Float32.init(inP as! Float16) + } else { + fatalError() + } + } + + public init(inFloat: Float32) { + self = inFloat + } + + public init(inFloat16: Float16) { + self = Float32.init(inFloat16) + } + + public static var bitSize: UInt { + return 32 + } } // 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 + + 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 } - set { - var newN = (Layout.H, newValue) - if let index = layoutWithDim.index(where: { (layout: Layout, dim: Int) -> Bool in - return layout == .H - }) { - fatalError() - } + } + 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 } - } - var W: Int? { - get { - for layoutDim in layoutWithDim { - if layoutDim.0 == .W { - return layoutDim.1 - } - } - return nil + } + 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 } - set { - var newN = (Layout.W, newValue) - if let index = layoutWithDim.index(where: { (layout: Layout, dim: Int) -> Bool in - return layout == .W - }) { - fatalError() - } + } + 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 } - } - - - 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] - } - } + } + 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.. Bool { + if lhs.layoutWithDim.count == rhs.layoutWithDim.count { + var result = true + for i in 0.. { - public let dim: [Int] - public let resultArr: [P] - public var intermediateResults: [Texture

]? - public let elapsedTime: Double - public init(inDim: [Int], inResult: [P], inElapsedTime: Double, inIntermediateResults: [Texture

]? = nil) { - dim = inDim - resultArr = inResult - elapsedTime = inElapsedTime - intermediateResults = inIntermediateResults - } + public let dim: [Int] + public let resultArr: [P] + public var intermediateResults: [Variant]? + public let elapsedTime: Double + public init(inDim: [Int], inResult: [P], inElapsedTime: Double, inIntermediateResults: [Variant]? = nil) { + dim = inDim + resultArr = inResult + elapsedTime = inElapsedTime + intermediateResults = inIntermediateResults + } } extension ResultHolder: CustomDebugStringConvertible, CustomStringConvertible { - public var debugDescription: String { - var str = "" - str += "Dim: \(dim) \n value:[ " - if resultArr.count < 20 { - for d in resultArr { - str += " \(d) " - } - } else { - for d in stride(from: 0, to: resultArr.count, by: resultArr.count/20) { - str += " \(resultArr[d]) " - } - } - str += " ]" - return str - } - - public var description: String { - return debugDescription + public var debugDescription: String { + var str = "" + str += "Dim: \(dim) \n value:[ " + if resultArr.count < 20 { + for d in resultArr { + str += " \(d) " + } + } else { + for d in stride(from: 0, to: resultArr.count, by: resultArr.count/20) { + str += " \(resultArr[d]) " + } } + str += " ]" + return str + } + + public var description: String { + return debugDescription + } } public class Executor { - var ops: [Runable & InferShaperable] = [] - let program: Program - let device: MTLDevice - let queue: MTLCommandQueue - public init(inDevice:MTLDevice, inQueue: MTLCommandQueue, inProgram: Program) throws { - program = inProgram - device = inDevice - queue = inQueue - for block in inProgram.programDesc.blocks { - //block.ops.count - for i in 0...shared.creat(device: inDevice, opDesc: op, scope: inProgram.scope) - op.inferShape() - ops.append(op) - } catch let error { - throw error - } - } + var ops: [Runable & InferShaperable] = [] + let program: Program + let device: MTLDevice + let queue: MTLCommandQueue + public init(inDevice:MTLDevice, inQueue: MTLCommandQueue, inProgram: Program) throws { + program = inProgram + device = inDevice + queue = inQueue + for block in inProgram.programDesc.blocks { + //block.ops.count + for i in 0..<39 { + let op = block.ops[i] + do { + let op = try OpCreator

.shared.creat(device: inDevice, opDesc: op, scope: inProgram.scope) +// op.inferShape() + ops.append(op) + } catch let error { + throw error } - } + } + } + } + + public func predict(input: MTLTexture, expect: [Int], completionHandle: @escaping (ResultHolder

) -> Void, preProcessKernle: CusomKernel? = nil, except: Int = 0) throws { + 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 + } - public func predict(input: MTLTexture, expect: [Int], completionHandle: @escaping (ResultHolder

) -> Void, preProcessKernle: CusomKernel? = nil) throws { - 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) - } catch let error { - throw error - } - } - - buffer.addCompletedHandler { (commandbuffer) in -// let inputArr = resInput.floatArray(res: { (p:P) -> P in -// return p -// }) -// print(inputArr) - -// let stridableInput: [(index: Int, value: Float)] = input.stridableFloatArray() -// print(stridableInput) - -// let _: Flo? = input.logDesc(header: "input: ", stridable: true) -// for op in self.ops { -// op.delogOutput() -// } -// return - -// self.ops[2].delogOutput() - - let afterDate = Date.init() - - guard let outputVar = self.program.scope.output() else { - fatalError("output nil") - } - - guard let output = outputVar as? Texture

else { - fatalError("output var type error") - } - let resultHodlder = ResultHolder

.init(inDim: output.dim.dims, inResult: output.metalTexture.floatArray(res: { (p:P) -> P in - return p - }), inElapsedTime: afterDate.timeIntervalSince(beforeDate)) - completionHandle(resultHodlder) - } - buffer.commit() + let beforeDate = Date.init() + let inputTexture = InputTexture.init(inMTLTexture: resInput, inExpectDim: Dim.init(inDim: expect)) + program.scope.setInput(input: inputTexture) + //(ops.count - except) + for i in 0.. 0 { + outputTextures = ops[ops.count - except].inputs() } + buffer.addCompletedHandler { (commandbuffer) in +// return; + +// let inputArr = resInput.floatArray(res: { (p:P) -> P in +// return p +// }) + +// writeToLibrary(fileName: "input_hand", array: inputArr) +// print("write to library done") +// return + // print(inputArr) + + // let stridableInput: [(index: Int, value: Float)] = input.stridableFloatArray() + // print(stridableInput) + + // let _: Flo? = input.logDesc(header: "input: ", stridable: true) + for op in self.ops { +// op.delogOutput() + } +// return + + self.ops[38].delogOutput() +// self.ops[91].delogOutput() +// self.ops[92].delogOutput() +// self.ops[93].delogOutput() + + return; + + let afterDate = Date.init() + + var resultHolder: ResultHolder

+ if except > 0 { + resultHolder = ResultHolder

.init(inDim: [], inResult: [], inElapsedTime: afterDate.timeIntervalSince(beforeDate), inIntermediateResults: outputTextures) + } else { + let outputVar: Variant = self.program.scope.output()! + let output: Texture

= outputVar as! Texture

+ + resultHolder = ResultHolder

.init(inDim: output.dim.dims, inResult: output.metalTexture.floatArray(res: { (p:P) -> P in + return p + }), inElapsedTime: afterDate.timeIntervalSince(beforeDate)) + } + + + completionHandle(resultHolder) + } + buffer.commit() + } + + public func clear() { + program.scope.clear() + } + } //public let paddle_executor: Executor = Executor.init() diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift b/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift index 952004a00d4fcac70c45164893a9b35f8d177b9a..44f7bfd20ff9aa9e3c9dce1da3850acc4994cd46 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift @@ -16,100 +16,101 @@ import Metal import Foundation protocol Fusion { - static func fusionNode() -> Node - static func change() -> [String : [(from: String, to: String)]] - static func fusionType() -> String + static func fusionNode() -> Node + static func change() -> [String : [(from: String, to: String)]] + static func fusionType() -> String } protocol Runable { - func run(device: MTLDevice, buffer: MTLCommandBuffer) throws - func runImpl(device: MTLDevice,buffer: MTLCommandBuffer) throws - func delogOutput() + func run(device: MTLDevice, buffer: MTLCommandBuffer) throws + func runImpl(device: MTLDevice,buffer: MTLCommandBuffer) throws + func delogOutput() + func inputs() -> [Variant] } extension Runable where Self: OperatorProtocol{ - func run(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try runImpl(device: device, buffer: buffer) - } catch let error { - throw error - } -// print(type + ": " + para.outputDesc()) - } - - func delogOutput() { - print(type + ": has no implementation" ) + func run(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try runImpl(device: device, buffer: buffer) + } catch let error { + throw error } +// print(type + ": " + para.outputDesc()) + } + + func delogOutput() { + print(type + ": has no implementation" ) + } } protocol Creator where Self: OperatorProtocol{ - associatedtype OpType: OperatorProtocol & Runable & InferShaperable - static func creat(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws -> OpType + associatedtype OpType: OperatorProtocol & Runable & InferShaperable + static func creat(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws -> OpType } extension Creator where Self: OperatorProtocol { - static func creat(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws -> OpType { - do { - return try OpType.provide(device:device, opDesc: opDesc, inScope: inScope) - } catch let error { - throw error - } + static func creat(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws -> OpType { + do { + return try OpType.provide(device:device, opDesc: opDesc, inScope: inScope) + } catch let error { + throw error } + } } protocol InferShaperable { - func inferShape() + func inferShape() } protocol OperatorProtocol { - associatedtype ParamType - associatedtype KerType: Computable where Self.KerType.ParamType == ParamType - var type: String { get } - var scope: Scope { get } - var inputs: [String : [String]] { get } - var paraInputs: [String : [String]] { get set } - var outpus: [String : [String]] { get } - var attrs: [String : Attr] { get } - var para: ParamType { get } - var kernel: KerType { get } - init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws + associatedtype ParamType + associatedtype KerType: Computable where Self.KerType.ParamType == ParamType + var type: String { get } + var scope: Scope { get } + var inputs: [String : [String]] { get } + var paraInputs: [String : [String]] { get set } + var outpus: [String : [String]] { get } + var attrs: [String : Attr] { get } + var para: ParamType { get } + var kernel: KerType { get } + init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws } extension OperatorProtocol { - static func provide(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws -> Self { - do { - return try Self.init(device: device, opDesc: opDesc, inScope: inScope) - } catch let error { - throw error - } + static func provide(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws -> Self { + do { + return try Self.init(device: device, opDesc: opDesc, inScope: inScope) + } catch let error { + throw error } + } } class Operator : OperatorProtocol where KernelType.ParamType == ParameterType { - typealias ParamType = ParameterType - typealias KerType = KernelType - let type: String - let inputs: [String : [String]] - var paraInputs: [String : [String]] - let outpus: [String : [String]] - let attrs: [String : Attr] - let para: ParamType - let scope: Scope - var kernel: KerType - required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { - type = opDesc.type - scope = inScope - inputs = opDesc.inputs - outpus = opDesc.outputs - attrs = opDesc.attrs - paraInputs = opDesc.paraInputs - do { - para = try ParamType.init(opDesc:opDesc, inScope: inScope) - } catch let error { - throw error - } - kernel = KernelType.init(device: device, param: para) + typealias ParamType = ParameterType + typealias KerType = KernelType + let type: String + let inputs: [String : [String]] + var paraInputs: [String : [String]] + let outpus: [String : [String]] + let attrs: [String : Attr] + let para: ParamType + let scope: Scope + var kernel: KerType + required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { + type = opDesc.type + scope = inScope + inputs = opDesc.inputs + outpus = opDesc.outputs + attrs = opDesc.attrs + paraInputs = opDesc.paraInputs + do { + para = try ParamType.init(opDesc:opDesc, inScope: inScope) + } catch let error { + throw error } + kernel = KernelType.init(device: device, param: para) + } } // op infos diff --git a/metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift index 3761dad60f0f8b20e3f95168445317a3e627ada9..1ed6fa7865e09437de047e01328f4f6747e6b1de 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift @@ -15,45 +15,50 @@ import Foundation class BatchNormParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - input = try BatchNormParam.inputX(inputs: opDesc.inputs, from: inScope) - output = try BatchNormParam.outputY(outputs: opDesc.outputs, from: inScope) - inputBias = try BatchNormParam.inputBiase(inputs: opDesc.paraInputs, from: inScope) - inputMean = try BatchNormParam.inputMean(inputs: opDesc.paraInputs, from: inScope) - inputScale = try BatchNormParam.inputScale(inputs: opDesc.paraInputs, from: inScope) - inputVariance = try BatchNormParam.inputVariance(inputs: opDesc.paraInputs, from: inScope) - epsilon = try BatchNormParam.getAttr(key: "epsilon", attrs: opDesc.attrs) - momentum = try BatchNormParam.getAttr(key: "momentum", attrs: opDesc.attrs) - is_test = try BatchNormParam.getAttr(key: "is_test", attrs: opDesc.attrs) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + input = try BatchNormParam.inputX(inputs: opDesc.inputs, from: inScope) + output = try BatchNormParam.outputY(outputs: opDesc.outputs, from: inScope) + inputBias = try BatchNormParam.inputBiase(inputs: opDesc.paraInputs, from: inScope) + inputMean = try BatchNormParam.inputMean(inputs: opDesc.paraInputs, from: inScope) + inputScale = try BatchNormParam.inputScale(inputs: opDesc.paraInputs, from: inScope) + inputVariance = try BatchNormParam.inputVariance(inputs: opDesc.paraInputs, from: inScope) + epsilon = try BatchNormParam.getAttr(key: "epsilon", attrs: opDesc.attrs) + momentum = try BatchNormParam.getAttr(key: "momentum", attrs: opDesc.attrs) + is_test = try BatchNormParam.getAttr(key: "is_test", attrs: opDesc.attrs) + } catch let error { + throw error } - let input: Texture

- var output: Texture

- let inputBias: Tensor - let inputMean: Tensor - let inputScale: Tensor - let inputVariance: Tensor - let epsilon: Float - let momentum: Float - let is_test: Bool + } + let input: Texture

+ var output: Texture

+ let inputBias: Tensor + let inputMean: Tensor + let inputScale: Tensor + let inputVariance: Tensor + let epsilon: Float + let momentum: Float + let is_test: Bool } class BatchNormOp: Operator, BatchNormParam

>, Runable, Creator, InferShaperable{ - func inferShape() { - para.output.dim = para.input.dim - } - typealias OpType = BatchNormOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } + + func inputs() -> [Variant] { + return [para.input, para.inputBias, para.inputMean, para.inputScale, para.inputVariance] + } + + func inferShape() { + para.output.dim = para.input.dim + } + typealias OpType = BatchNormOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift index f50664dd1a7c165c5df06965cdca992492a86650..994997d491baab83dfbbc7348fb4c7915110b518 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift @@ -1,37 +1,37 @@ -///* 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 Foundation - -class BoxcoderParam: OpParam { + ///* 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 Foundation + + class BoxcoderParam: OpParam { typealias ParamPrecisionType = P required init(opDesc: OpDesc, inScope: Scope) throws { - do { - priorBox = try BoxcoderParam.getFirstTensor(key: "PriorBox", map: opDesc.inputs, from: inScope) - priorBoxVar = try BoxcoderParam.getFirstTensor(key: "PriorBoxVar", map: opDesc.inputs, from: inScope) - targetBox = try BoxcoderParam.getFirstTensor(key: "TargetBox", map: opDesc.inputs, from: inScope) - output = try BoxcoderParam.getFirstTensor(key: "OutputBox", map: opDesc.outputs, from: inScope) - codeType = try BoxcoderParam.getAttr(key: "code_type", attrs: opDesc.attrs) - boxNormalized = try BoxcoderParam.getAttr(key: "box_normalized", attrs: opDesc.attrs) - } catch let error { - throw error - } - assert(priorBox.transpose == [0, 1, 2, 3]) - assert(priorBoxVar.transpose == [0, 1, 2, 3]) - assert(targetBox.transpose == [0, 1, 2, 3]) - assert(codeType == "decode_center_size") // encode_center_size is not implemented - assert((targetBox.tensorDim.cout() == 3) && (targetBox.tensorDim[0] == 1)) // N must be 1 (only handle batch size = 1) + do { + priorBox = try BoxcoderParam.getFirstTensor(key: "PriorBox", map: opDesc.inputs, from: inScope) + priorBoxVar = try BoxcoderParam.getFirstTensor(key: "PriorBoxVar", map: opDesc.inputs, from: inScope) + targetBox = try BoxcoderParam.getFirstTensor(key: "TargetBox", map: opDesc.inputs, from: inScope) + output = try BoxcoderParam.getFirstTensor(key: "OutputBox", map: opDesc.outputs, from: inScope) + codeType = try BoxcoderParam.getAttr(key: "code_type", attrs: opDesc.attrs) + boxNormalized = try BoxcoderParam.getAttr(key: "box_normalized", attrs: opDesc.attrs) + } catch let error { + throw error + } + assert(priorBox.transpose == [0, 1, 2, 3]) + assert(priorBoxVar.transpose == [0, 1, 2, 3]) + assert(targetBox.transpose == [0, 1, 2, 3]) + assert(codeType == "decode_center_size") // encode_center_size is not implemented + assert((targetBox.tensorDim.cout() == 3) && (targetBox.tensorDim[0] == 1)) // N must be 1 (only handle batch size = 1) } let priorBox: Texture

let priorBoxVar: Texture

@@ -39,23 +39,42 @@ class BoxcoderParam: OpParam { var output: Texture

let codeType: String let boxNormalized: Bool -} - -class BoxcoderOp: Operator, BoxcoderParam

>, Runable, Creator, InferShaperable{ + } + + class BoxcoderOp: Operator, BoxcoderParam

>, Runable, Creator, InferShaperable{ + + func inputs() -> [Variant] { + return [para.priorBox, para.priorBoxVar, para.targetBox] + } func inferShape() { -// para.output.dim = para.input.dim + // para.output.dim = para.input.dim } typealias OpType = BoxcoderOp

func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } } -} - - - + + func delogOutput() { + let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in + return o + } + print(outputArray.strideArray()) + //box_coder_0.tmp_0 +// writeToLibrary(fileName: "boxcoder_output", array: outputArray) + print(para.output.metalTexture) + print(" write done ") + } + + } + + + + + + diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift index 1dd0a06ce7ee7bfee1bfcdefd74f772f3e5ce1cb..5f0004a726fdf11439d96628bb616826324c0a89 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift @@ -15,44 +15,67 @@ import Foundation class ConcatParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - guard let xlist = opDesc.inputs["X"] else { - fatalError() - } - for x in xlist { - guard let variant = inScope[x], let v = variant as? Texture

else { - fatalError() - } - input.append(v) - } - axis = try ConcatParam.getAttr(key: "axis", attrs: opDesc.attrs) - output = try ConcatParam.outputOut(outputs: opDesc.outputs, from: inScope) - } catch let error { - throw error + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + guard let xlist = opDesc.inputs["X"] else { + fatalError() + } + for x in xlist { + guard let variant = inScope[x], let v = variant as? Texture

else { + fatalError() } + input.append(v) + } + axis = try ConcatParam.getAttr(key: "axis", attrs: opDesc.attrs) + output = try ConcatParam.outputOut(outputs: opDesc.outputs, from: inScope) + } catch let error { + throw error } - var input: [Texture

] = [] - var output: Texture

- let axis: Int + } + var input: [Texture

] = [] + var output: Texture

+ let axis: Int } class ConcatOp: Operator, ConcatParam

>, Runable, Creator, InferShaperable{ - - func inferShape() { -// let dim = para.input.reduce([0, 0]) {[$0[0] + $1.dim[0], $1.dim[1]]} -// para.output.dim = Dim.init(inDim: dim) + + func inputs() -> [Variant] { + return para.input + } + + func inferShape() { + // let dim = para.input.reduce([0, 0]) {[$0[0] + $1.dim[0], $1.dim[1]]} + // para.output.dim = Dim.init(inDim: dim) + } + + typealias OpType = ConcatOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } - - typealias OpType = ConcatOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } + } + + func delogOutput() { + let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in + return o } + print(outputArray.strideArray()) + let device: MTLDevice = MTLCreateSystemDefaultDevice()! + +// let tensorArray: [P] = device.texture2tensor(texture: para.output.metalTexture, dim: [1917, 4]) + +// print(tensorArray.strideArray()) + +// print(para.output.metalTexture) + +// writeToLibrary(fileName: "concat_out", array: outputArray) +// print(" write done ") + +// print(outputArray.strideArray()) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift index ce47ac434237b463f52e7f84a335a549553082c2..f037cc69287022a9be39fc0416daed0001dda86a 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddBatchNormReluOp.swift @@ -16,120 +16,125 @@ import Foundation class ConvAddBatchNormReluParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - - filter = try ConvAddBatchNormReluParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) - input = try ConvAddBatchNormReluParam.input(inputs: opDesc.inputs, from: inScope) - output = try ConvAddBatchNormReluParam.outputOut(outputs: opDesc.outputs, from: inScope) - stride = try ConvAddBatchNormReluParam.getAttr(key: "strides", attrs: opDesc.attrs) - paddings = try ConvAddBatchNormReluParam.getAttr(key: "paddings", attrs: opDesc.attrs) - dilations = try ConvAddBatchNormReluParam.getAttr(key: "dilations", attrs: opDesc.attrs) - epsilon = try ConvAddBatchNormReluParam.getAttr(key: "epsilon", attrs: opDesc.attrs) - - groups = try ConvAddBatchNormReluParam.getAttr(key: "groups", attrs: opDesc.attrs) - variance = try ConvAddBatchNormReluParam.inputVariance(inputs: opDesc.paraInputs, from: inScope) - bias = try ConvAddBatchNormReluParam.inputBiase(inputs: opDesc.paraInputs, from: inScope) - - scale = try ConvAddBatchNormReluParam.inputScale(inputs: opDesc.paraInputs, from: inScope) - mean = try ConvAddBatchNormReluParam.inputMean(inputs: opDesc.paraInputs, from: inScope) - y = try ConvAddBatchNormReluParam.inputY(inputs: opDesc.paraInputs, from: inScope) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + + filter = try ConvAddBatchNormReluParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) + input = try ConvAddBatchNormReluParam.input(inputs: opDesc.inputs, from: inScope) + output = try ConvAddBatchNormReluParam.outputOut(outputs: opDesc.outputs, from: inScope) + stride = try ConvAddBatchNormReluParam.getAttr(key: "strides", attrs: opDesc.attrs) + paddings = try ConvAddBatchNormReluParam.getAttr(key: "paddings", attrs: opDesc.attrs) + dilations = try ConvAddBatchNormReluParam.getAttr(key: "dilations", attrs: opDesc.attrs) + epsilon = try ConvAddBatchNormReluParam.getAttr(key: "epsilon", attrs: opDesc.attrs) + + groups = try ConvAddBatchNormReluParam.getAttr(key: "groups", attrs: opDesc.attrs) + variance = try ConvAddBatchNormReluParam.inputVariance(inputs: opDesc.paraInputs, from: inScope) + bias = try ConvAddBatchNormReluParam.inputBiase(inputs: opDesc.paraInputs, from: inScope) + + scale = try ConvAddBatchNormReluParam.inputScale(inputs: opDesc.paraInputs, from: inScope) + mean = try ConvAddBatchNormReluParam.inputMean(inputs: opDesc.paraInputs, from: inScope) + y = try ConvAddBatchNormReluParam.inputY(inputs: opDesc.paraInputs, from: inScope) + } catch let error { + throw error } - - let input: Texture

- - let variance: Tensor - let bias: Tensor - let mean: Tensor - let scale: Tensor - let y: Tensor - let filter: Tensor - let epsilon: Float32 - var newScale: MTLBuffer? - var newBiase: MTLBuffer? - - var output: Texture

- let stride: [Int32] - let paddings: [Int32] - let dilations: [Int32] - let groups: Int + } + + let input: Texture

+ + let variance: Tensor + let bias: Tensor + let mean: Tensor + let scale: Tensor + let y: Tensor + let filter: Tensor + let epsilon: Float32 + var newScale: MTLBuffer? + var newBiase: MTLBuffer? + + var output: Texture

+ let stride: [Int32] + let paddings: [Int32] + let dilations: [Int32] + let groups: Int } class ConvAddBatchNormReluOp: Operator, ConvAddBatchNormReluParam

>, Runable, Creator, InferShaperable, Fusion{ - typealias OpType = ConvAddBatchNormReluOp

+ + func inputs() -> [Variant] { + return [para.variance, para.bias, para.mean, para.scale, para.y, para.filter, para.input] + } + + typealias OpType = ConvAddBatchNormReluOp

+ + func inferShape() { + let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations - func inferShape() { - let inDims = para.input.dim - let filterDim = para.filter.dim - let strides = para.stride - let paddings = para.paddings - let dilations = para.dilations - - var outDim = [inDims[0]] - for i in 0.. Node { + let beginNode = Node.init(inType: gConvType) + _ = beginNode + --> Node.init(inType: gElementwiseAddType) + --> Node.init(inType: gBatchNormType) + --> Node.init(inType: gReluType) + return beginNode + } + + static func change() -> [String : [(from: String, to: String)]] { + return [:] + } + + static func fusionType() -> String { + return gConvAddBatchNormReluType + } + + func delogOutput() { - static func fusionNode() -> Node { - let beginNode = Node.init(inType: gConvType) - _ = beginNode - --> Node.init(inType: gElementwiseAddType) - --> Node.init(inType: gBatchNormType) - --> Node.init(inType: gReluType) - return beginNode - } + // 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)") - static func change() -> [String : [(from: String, to: String)]] { - return [:] - } + // print("biase: \(para.y)") + // print("padding: \(para.paddings)") + // print("stride: \(para.stride)") - static func fusionType() -> String { - return gConvAddBatchNormReluType - } + // let _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false) + // let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false) + // let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false) - func delogOutput() { - -// 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.y)") -// print("padding: \(para.paddings)") -// print("stride: \(para.stride)") - -// let _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false) -// let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false) -// let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false) - - let output = para.output.metalTexture.floatArray { (p: P) -> P in - return p - } -// - writeToLibrary(fileName: "output_112x112x32_2", array: output) - print(" write done") - -// let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false) + let output = para.output.metalTexture.floatArray { (p: P) -> P in + return p } + // + writeToLibrary(fileName: "output_112x112x32_2", array: output) + print(" write done") + + // let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift index f1ab27bc0de2e08c409794e813c86caa0f18ceb6..60475be70159c5947b4fef7198bd0dd8c9c4d26f 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvAddOp.swift @@ -15,80 +15,108 @@ import Foundation class ConvAddParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - filter = try ConvAddParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) - input = try ConvAddParam.input(inputs: opDesc.inputs, from: inScope) - output = try ConvAddParam.outputOut(outputs: opDesc.outputs, from: inScope) - stride = try ConvAddParam.getAttr(key: "strides", attrs: opDesc.attrs) - paddings = try ConvAddParam.getAttr(key: "paddings", attrs: opDesc.attrs) - dilations = try ConvAddParam.getAttr(key: "dilations", attrs: opDesc.attrs) - groups = try ConvAddParam.getAttr(key: "groups", attrs: opDesc.attrs) - y = try ConvAddParam.inputY(inputs: opDesc.paraInputs, from: inScope) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + filter = try ConvAddParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) + input = try ConvAddParam.input(inputs: opDesc.inputs, from: inScope) + output = try ConvAddParam.outputOut(outputs: opDesc.outputs, from: inScope) + stride = try ConvAddParam.getAttr(key: "strides", attrs: opDesc.attrs) + paddings = try ConvAddParam.getAttr(key: "paddings", attrs: opDesc.attrs) + dilations = try ConvAddParam.getAttr(key: "dilations", attrs: opDesc.attrs) + groups = try ConvAddParam.getAttr(key: "groups", attrs: opDesc.attrs) + y = try ConvAddParam.inputY(inputs: opDesc.paraInputs, from: inScope) + } catch let error { + throw error } - - let input: Texture

- let y: Tensor - let filter: Tensor - - var output: Texture

- let stride: [Int32] - let paddings: [Int32] - let dilations: [Int32] - let groups: Int + } + + let input: Texture

+ let y: Tensor + let filter: Tensor + + var output: Texture

+ let stride: [Int32] + let paddings: [Int32] + let dilations: [Int32] + let groups: Int } class ConvAddOp: Operator, ConvAddParam

>, Runable, Creator, InferShaperable, Fusion{ - static func fusionNode() -> Node { - let beginNode = Node.init(inType: gConvType) - _ = beginNode - --> Node.init(inType: gElementwiseAddType) - return beginNode - } + + func delogOutput() { + print(" conv add: ") + +// print(para.input.metalTexture) - static func change() -> [String : [(from: String, to: String)]] { - return [:] +// print(" filter array: ") +// let filterArray: [P] = para.filter.buffer.array() +// print(filterArray) + + let input = para.input.metalTexture.floatArray { (p: P) -> P in + return p } +// print(input) - static func fusionType() -> String { - return gConvAddType + let output = para.output.metalTexture.floatArray { (p: P) -> P in + return p } +// print(para.output.metalTexture) + print(output) + } + + + static func fusionNode() -> Node { + let beginNode = Node.init(inType: gConvType) + _ = beginNode + --> Node.init(inType: gElementwiseAddType) + return beginNode + } + + static func change() -> [String : [(from: String, to: String)]] { + return [:] + } + + func inputs() -> [Variant] { + return [para.input, para.y, para.filter] + } + + + static func fusionType() -> String { + return gConvAddType + } + + typealias OpType = ConvAddOp

+ + func inferShape() { - typealias OpType = ConvAddOp

+ let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations - func inferShape() { - - let inDims = para.input.dim - let filterDim = para.filter.dim - let strides = para.stride - let paddings = para.paddings - let dilations = para.dilations - - var outDim = [inDims[0]] - for i in 0..: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - filter = try ConvBNReluParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) - input = try ConvBNReluParam.input(inputs: opDesc.inputs, from: inScope) - output = try ConvBNReluParam.outputOut(outputs: opDesc.outputs, from: inScope) - stride = try ConvBNReluParam.getAttr(key: "strides", attrs: opDesc.attrs) - paddings = try ConvBNReluParam.getAttr(key: "paddings", attrs: opDesc.attrs) - dilations = try ConvBNReluParam.getAttr(key: "dilations", attrs: opDesc.attrs) - epsilon = try ConvBNReluParam.getAttr(key: "epsilon", attrs: opDesc.attrs) - - groups = try ConvBNReluParam.getAttr(key: "groups", attrs: opDesc.attrs) - variance = try ConvBNReluParam.inputVariance(inputs: opDesc.paraInputs, from: inScope) - bias = try ConvBNReluParam.inputBiase(inputs: opDesc.paraInputs, from: inScope) - scale = try ConvBNReluParam.inputScale(inputs: opDesc.paraInputs, from: inScope) - mean = try ConvBNReluParam.inputMean(inputs: opDesc.paraInputs, from: inScope) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + filter = try ConvBNReluParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) + input = try ConvBNReluParam.input(inputs: opDesc.inputs, from: inScope) + output = try ConvBNReluParam.outputOut(outputs: opDesc.outputs, from: inScope) + stride = try ConvBNReluParam.getAttr(key: "strides", attrs: opDesc.attrs) + paddings = try ConvBNReluParam.getAttr(key: "paddings", attrs: opDesc.attrs) + dilations = try ConvBNReluParam.getAttr(key: "dilations", attrs: opDesc.attrs) + epsilon = try ConvBNReluParam.getAttr(key: "epsilon", attrs: opDesc.attrs) + + groups = try ConvBNReluParam.getAttr(key: "groups", attrs: opDesc.attrs) + variance = try ConvBNReluParam.inputVariance(inputs: opDesc.paraInputs, from: inScope) + bias = try ConvBNReluParam.inputBiase(inputs: opDesc.paraInputs, from: inScope) + scale = try ConvBNReluParam.inputScale(inputs: opDesc.paraInputs, from: inScope) + mean = try ConvBNReluParam.inputMean(inputs: opDesc.paraInputs, from: inScope) + } catch let error { + throw error } - - let input: Texture

- - let variance: Tensor - let bias: Tensor - let mean: Tensor - let scale: Tensor - let filter: Tensor - let epsilon: Float32 - var newScale: MTLBuffer? - var newBiase: MTLBuffer? - - var output: Texture

- let stride: [Int32] - let paddings: [Int32] - let dilations: [Int32] - let groups: Int + } + + let input: Texture

+ + let variance: Tensor + let bias: Tensor + let mean: Tensor + let scale: Tensor + let filter: Tensor + let epsilon: Float32 + var newScale: MTLBuffer? + var newBiase: MTLBuffer? + + var output: Texture

+ let stride: [Int32] + let paddings: [Int32] + let dilations: [Int32] + let groups: Int } class ConvBNReluOp: Operator, ConvBNReluParam

>, Runable, Creator, InferShaperable, Fusion{ - typealias OpType = ConvBNReluOp

+ typealias OpType = ConvBNReluOp

+ + func inputs() -> [Variant] { + return [para.input, para.variance, para.bias, para.mean, para.scale, para.filter] + } + + + func inferShape() { + let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations - func inferShape() { - let inDims = para.input.dim - let filterDim = para.filter.dim - let strides = para.stride - let paddings = para.paddings - let dilations = para.dilations - - var outDim = [inDims[0]] - for i in 0.. Node { + let beginNode = Node.init(inType: gConvType) + _ = beginNode + --> Node.init(inType: gBatchNormType) + --> Node.init(inType: gReluType) + return beginNode + } + + static func change() -> [String : [(from: String, to: String)]] { + return [:] + } + + static func fusionType() -> String { + return gConvBnReluType + } + + func delogOutput() { - static func fusionNode() -> Node { - let beginNode = Node.init(inType: gConvType) - _ = beginNode - --> Node.init(inType: gBatchNormType) - --> Node.init(inType: gReluType) - return beginNode - } + // 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)") - static func change() -> [String : [(from: String, to: String)]] { - return [:] - } + // print("biase: \(para.y)") + // print("padding: \(para.paddings)") + // print("stride: \(para.stride)") + + // let _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false) + // let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false) + // let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false) + - static func fusionType() -> String { - return gConvBnReluType - } - func delogOutput() { - - // 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.y)") - // print("padding: \(para.paddings)") - // print("stride: \(para.stride)") - - // let _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false) - // let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false) - // let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false) - - let output = para.output.metalTexture.floatArray { (p: P) -> P in - return p - } - // - writeToLibrary(fileName: "output_112x112x32_2", array: output) - print(" write done") - - // let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false) +// print("input: ") +// print(para.input.metalTexture) +// +// let input = para.input.metalTexture.floatArray { (p: P) -> P in +// return p +// } +// for i in 0...size) { +// print("index: \(i) \(newScale![i]) ") +// } +// +// print("new biase: ") +// for i in 0..<(para.newBiase!.length / MemoryLayout

.size) { +// print("index: \(i) \(newBiase![i]) ") +// } + + print(para.output.metalTexture) + + let output = para.output.metalTexture.floatArray { (p: P) -> P in + return p } + print(output) +// + writeToLibrary(fileName: "batch_norm_34.tmp_2", array: output) + print(" write done") +// + + + + + // let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: true) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift index 29b0c4246e728dbc3d3b865a189c7063ac1bbdcf..8a76a39284bba84cd87b7df12f4516efc778fdc7 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift @@ -15,74 +15,79 @@ import Foundation class ConvParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - filter = try ConvParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) - input = try ConvParam.input(inputs: opDesc.inputs, from: inScope) - output = try ConvParam.output(outputs: opDesc.outputs, from: inScope) - stride = try ConvParam.getAttr(key: "strides", attrs: opDesc.attrs) - paddings = try ConvParam.getAttr(key: "paddings", attrs: opDesc.attrs) - dilations = try ConvParam.getAttr(key: "dilations", attrs: opDesc.attrs) - groups = try ConvParam.getAttr(key: "groups", attrs: opDesc.attrs) - - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + filter = try ConvParam.inputFilter(paraInputs: opDesc.paraInputs, from: inScope) + input = try ConvParam.input(inputs: opDesc.inputs, from: inScope) + output = try ConvParam.output(outputs: opDesc.outputs, from: inScope) + stride = try ConvParam.getAttr(key: "strides", attrs: opDesc.attrs) + paddings = try ConvParam.getAttr(key: "paddings", attrs: opDesc.attrs) + dilations = try ConvParam.getAttr(key: "dilations", attrs: opDesc.attrs) + groups = try ConvParam.getAttr(key: "groups", attrs: opDesc.attrs) + + } catch let error { + throw error } - - let input: Texture

- let filter: Tensor - var output: Texture

- let stride: [Int32] - let paddings: [Int32] - let dilations: [Int32] - let groups: Int + } + + let input: Texture

+ let filter: Tensor + var output: Texture

+ let stride: [Int32] + let paddings: [Int32] + let dilations: [Int32] + let groups: Int } class ConvOp: Operator, ConvParam

>, Runable, Creator, InferShaperable { - required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { - do { - try super.init(device: device, opDesc: opDesc, inScope: inScope) - } catch let error { - throw error - } - - } - func inferShape() { - let inDims = para.input.dim - let filterDim = para.filter.dim - let strides = para.stride - let paddings = para.paddings - let dilations = para.dilations - - var outDim = [inDims[0]] - for i in 0.. [Variant] { + return [para.input, para.filter] + } + + required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { + do { + try super.init(device: device, opDesc: opDesc, inScope: inScope) + } catch let error { + throw error } - typealias OpType = ConvOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } - } + } + func inferShape() { + let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations - func delogOutput() { - print("conv output : ") - print(para.output.metalTexture) -// let _: Float16? = para.output.metalTexture.logDesc() + var outDim = [inDims[0]] + for i in 0.. + func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } + } + + func delogOutput() { + print("conv output : ") + print(para.output.metalTexture) + // let _: Float16? = para.output.metalTexture.logDesc() + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvTransposeOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvTransposeOp.swift index 10dbecb7a46f150528df19a9ae1c3069d95ddc6a..206828e3a0a631061d60ac480dcf870f32f9737f 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConvTransposeOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvTransposeOp.swift @@ -28,6 +28,10 @@ class ConvTransposeParam: ConvParam

{ class ConvTransposeOp: Operator, ConvTransposeParam

>, Runable, Creator, InferShaperable{ + func inputs() -> [Variant] { + return [para.input, para.filter] + } + func inferShape() { // para.output.dim = para.input.dim } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/DepthwiseConvOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/DepthwiseConvOp.swift index 272c386cddd52b95f61c2e9aca3aa72cd997cc1e..b6f1e116b1f848857a1b04d038e34af416d98dd5 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/DepthwiseConvOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/DepthwiseConvOp.swift @@ -15,49 +15,54 @@ import Foundation class DepthConvOp: Operator, ConvParam

>, Runable, Creator, InferShaperable { - required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { - do { - try super.init(device: device, opDesc: opDesc, inScope: inScope) - } catch let error { - throw error - } + + func inputs() -> [Variant] { + return [para.input, para.filter] + } + + required init(device: MTLDevice, opDesc: OpDesc, inScope: Scope) throws { + do { + try super.init(device: device, opDesc: opDesc, inScope: inScope) + } catch let error { + throw error } + } + + func inferShape() { + let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations - func inferShape() { - let inDims = para.input.dim - let filterDim = para.filter.dim - let strides = para.stride - let paddings = para.paddings - let dilations = para.dilations - - var outDim = [inDims[0]] - for i in 0.. - - func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } - } - - func delogOutput() { - print("conv output : ") - print(para.output.metalTexture) - // let _: Float16? = para.output.metalTexture.logDesc() + outDim.append(filterDim[0]) + para.output.dim = Dim.init(inDim: outDim) + } + + typealias OpType = DepthConvOp

+ + func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } + } + + func delogOutput() { + print("conv output : ") + print(para.output.metalTexture) + // let _: Float16? = para.output.metalTexture.logDesc() + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/DwConvBNReluOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/DwConvBNReluOp.swift index da6651d94ff299cbe91aa547b3ae51d68c0aa4ec..15d2b262f332fc9fb7af6da611c05ff9ac4b0249 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/DwConvBNReluOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/DwConvBNReluOp.swift @@ -15,75 +15,79 @@ import Foundation class DwConvBNReluOp: Operator, ConvBNReluParam

>, Runable, Creator, InferShaperable, Fusion{ - typealias OpType = ConvBNReluOp

+ typealias OpType = ConvBNReluOp

+ + func inputs() -> [Variant] { + return [para.input, para.bias, para.mean, para.filter, para.variance, para.scale] + } + + func inferShape() { + let inDims = para.input.dim + let filterDim = para.filter.dim + let strides = para.stride + let paddings = para.paddings + let dilations = para.dilations - func inferShape() { - let inDims = para.input.dim - let filterDim = para.filter.dim - let strides = para.stride - let paddings = para.paddings - let dilations = para.dilations - - var outDim = [inDims[0]] - for i in 0.. Node { + let beginNode = Node.init(inType: gDepthConvType) + _ = beginNode + --> Node.init(inType: gBatchNormType) + --> Node.init(inType: gReluType) + return beginNode + } + + static func change() -> [String : [(from: String, to: String)]] { + return [:] + } + + static func fusionType() -> String { + return gDwConvBnReluType + } + + func delogOutput() { - static func fusionNode() -> Node { - let beginNode = Node.init(inType: gDepthConvType) - _ = beginNode - --> Node.init(inType: gBatchNormType) - --> Node.init(inType: gReluType) - return beginNode - } + // 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)") - static func change() -> [String : [(from: String, to: String)]] { - return [:] - } + // print("biase: \(para.y)") + // print("padding: \(para.paddings)") + // print("stride: \(para.stride)") - static func fusionType() -> String { - return gDwConvBnReluType - } + // let _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false) + // let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false) + // let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false) - func delogOutput() { - - // 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.y)") - // print("padding: \(para.paddings)") - // print("stride: \(para.stride)") - - // let _: P? = para.y.buffer?.logDesc(header: " biase: ", stridable: false) - // let _: P? = para.newBiase?.logDesc(header: "new biase: ", stridable: false) - // let _: P? = para.newScale?.logDesc(header: "new scale: ", stridable: false) - - let output = para.output.metalTexture.floatArray { (p: P) -> P in - return p - } - // - writeToLibrary(fileName: "output_112x112x32_2", array: output) - print(" write done") - - // let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false) - } +// let output = para.output.metalTexture.floatArray { (p: P) -> P in +// return p +// } +// +// writeToLibrary(fileName: "batch_norm_19.tmp_2", array: output) +// print(" write done") + + // let _: P? = para.output.metalTexture.logDesc(header: "conv add batchnorm relu output: ", stridable: false) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ElementwiseAddOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ElementwiseAddOp.swift index 5ed36f86d79ffd639dc2ba76da74d24a532b1bd1..b4760f1afe489ec13238aeb6d50ac1e416d88478 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ElementwiseAddOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ElementwiseAddOp.swift @@ -15,33 +15,37 @@ import Foundation class ElementwiseAddParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - input = try ElementwiseAddParam.inputX(inputs: opDesc.inputs, from: inScope) - inputY = try ElementwiseAddParam.inputY(inputs: opDesc.paraInputs, from: inScope) - - output = try ElementwiseAddParam.outputOut(outputs: opDesc.outputs, from: inScope) - axis = try ElementwiseAddParam.getAttr(key: "axis", attrs: opDesc.attrs) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + input = try ElementwiseAddParam.inputX(inputs: opDesc.inputs, from: inScope) + inputY = try ElementwiseAddParam.inputY(inputs: opDesc.paraInputs, from: inScope) + + output = try ElementwiseAddParam.outputOut(outputs: opDesc.outputs, from: inScope) + axis = try ElementwiseAddParam.getAttr(key: "axis", attrs: opDesc.attrs) + } catch let error { + throw error } - let input: Texture

- let inputY: Tensor

- var output: Texture

- let axis: Int + } + let input: Texture

+ let inputY: Tensor

+ var output: Texture

+ let axis: Int } class ElementwiseAddOp: Operator, ElementwiseAddParam

>, Runable, Creator, InferShaperable{ - - func inferShape() { - para.output.dim = para.input.dim - } - - typealias OpType = ElementwiseAddOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - } + + func inputs() -> [Variant] { + return [para.input, para.inputY] + } + + func inferShape() { + para.output.dim = para.input.dim + } + + typealias OpType = ElementwiseAddOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift index c81d9e786c91408d2412b30eaec089904df75751..43554a84d6fc1a0d946c6c8da9a5bb10cf2e52e7 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift @@ -15,54 +15,58 @@ import Foundation class FeedParam: OpParam{ - var output: Texture

- var input: InputTexture { - return scope.input() as! InputTexture + var output: Texture

+ var input: InputTexture { + return scope.input() as! InputTexture + } + let scope: Scope + + required init(opDesc: OpDesc, inScope: Scope) throws { + scope = inScope + do { + output = try FeedParam.outputOut(outputs: opDesc.outputs, from: inScope) + } catch let error { + throw error } - let scope: Scope - - required init(opDesc: OpDesc, inScope: Scope) throws { - scope = inScope - do { - output = try FeedParam.outputOut(outputs: opDesc.outputs, from: inScope) - } catch let error { - throw error - } - } - - typealias ParamPrecisionType = P + } + + typealias ParamPrecisionType = P } class FeedOp: Operator, FeedParam

>, Runable, Creator, InferShaperable { - typealias OpType = FeedOp

- - func inferShape() { - // print("feed input: \(para.input.expectDim)") - print("feed output: \(para.output.dim)") - // para.output.dim = - // para.output.dim = para.input.expectDim + typealias OpType = FeedOp

+ + func inputs() -> [Variant] { + return [para.input] + } + + func inferShape() { + // print("feed input: \(para.input.expectDim)") + print("feed output: \(para.output.dim)") + // para.output.dim = + // para.output.dim = para.input.expectDim + } + + func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } - func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } - -// let resizeKernel = ResizeKernel

.init(device: device) -// let resizeParam = ResizeParam.init(input: para.input.mtlTexture, output: para.output.metalTexture, expectDim: para.input.expectDim) -// do { -// try resizeKernel.compute(commandBuffer: buffer, param: resizeParam) -// } catch let error { -// throw error -// } - } - - func delogOutput() { -// para.input.mtlTexture.logDesc() -// let _: P? = para.input.mtlTexture.logDesc(header: "feed input: ", stridable: true) -// let _: P? = para.output.metalTexture.logDesc(header: "feed output: ", stridable: false) - } + // let resizeKernel = ResizeKernel

.init(device: device) + // let resizeParam = ResizeParam.init(input: para.input.mtlTexture, output: para.output.metalTexture, expectDim: para.input.expectDim) + // do { + // try resizeKernel.compute(commandBuffer: buffer, param: resizeParam) + // } catch let error { + // throw error + // } + } + + func delogOutput() { + // para.input.mtlTexture.logDesc() + // let _: P? = para.input.mtlTexture.logDesc(header: "feed input: ", stridable: true) + // let _: P? = para.output.metalTexture.logDesc(header: "feed output: ", stridable: false) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift index 2964b89e5ddabbbbd4f2df032efa5ef2db82ec96..1bedd030b1232d4086a5d67455a57c940ed0e1a1 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift @@ -15,40 +15,44 @@ import Foundation class FetchParam: OpParam{ - var output: Texture

- let input: Texture

- let scope: Scope - required init(opDesc: OpDesc, inScope: Scope) throws { - scope = inScope - do { - input = try FetchParam.inputX(inputs: opDesc.inputs, from: inScope) - output = input - } catch let error { - throw error - } + var output: Texture

+ let input: Texture

+ let scope: Scope + required init(opDesc: OpDesc, inScope: Scope) throws { + scope = inScope + do { + input = try FetchParam.inputX(inputs: opDesc.inputs, from: inScope) + output = input + } catch let error { + throw error } - - typealias ParamPrecisionType = P + } + + typealias ParamPrecisionType = P } class FetchKernel: Kernel, Computable { - - func compute(commandBuffer: MTLCommandBuffer, param: FetchParam

) throws { - } - - required init(device: MTLDevice, param: FetchParam

) { - super.init(device: device, inFunctionName: "texture2d_to_2d_array") - } + + func compute(commandBuffer: MTLCommandBuffer, param: FetchParam

) throws { + } + + required init(device: MTLDevice, param: FetchParam

) { + super.init(device: device, inFunctionName: "texture2d_to_2d_array") + } } class FetchOp: Operator< FetchKernel

, FetchParam

>, Runable, Creator, InferShaperable{ - func inferShape() { - print(para.input.dim) - } - - typealias OpType = FetchOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - scope.setOutput(output: para.output) - } + func inputs() -> [Variant] { + return [para.input] + } + + func inferShape() { + print(para.input.dim) + } + + typealias OpType = FetchOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + scope.setOutput(output: para.output) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift index 88e5f015684221bd7286e0ee4e608d85acaa325d..f2bc495158ef14ab60647ebae34c78eb6802dd56 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddBatchNormReluKernel.swift @@ -50,7 +50,7 @@ class ConvAddBatchNormReluKernel: Kernel, Computable, Testable required init(device: MTLDevice, param: ConvAddBatchNormReluParam

) { - param.output.initTexture(device: device, transpose: [0, 2, 3, 1]) + param.output.initTexture(device: device, inTranspose: [0, 2, 3, 1]) if param.filter.width == 1 && param.filter.height == 1 { super.init(device: device, inFunctionName: "conv_add_batch_norm_relu_1x1") diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift index d37100b27486adf1ccff6904e355ec3a66abdddb..033c0b7a1341ec86246b931ac0c250f6e719351f 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvAddKernel.swift @@ -25,7 +25,7 @@ class ConvAddKernel: Kernel, Computable { super.init(device: device, inFunctionName: "conv_add_3x3") } - param.output.initTexture(device: device, transpose: [0, 3, 1, 2]) + param.output.initTexture(device: device, inTranspose: [0, 3, 2, 1]) 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/ConvBNReluKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift index de77121c1a92b8ec13dade910ab0c37d74048d13..2878d8287912c302e634398d6ed7609a012c9b36 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift @@ -59,7 +59,7 @@ class ConvBNReluKernel: Kernel, Computable, Testable { } else { super.init(device: device, inFunctionName: "conv_batch_norm_relu_3x3") } - param.output.initTexture(device: device, transpose: [0, 2, 3, 1]) + param.output.initTexture(device: device, inTranspose: [0, 2, 3, 1]) param.filter.initBuffer(device: device, precision: Tensor.BufferPrecision.Float32) param.variance.initBuffer(device: device) @@ -70,8 +70,13 @@ class ConvBNReluKernel: Kernel, Computable, Testable { let offsetX = param.filter.width/2 - Int(param.paddings[0]) let offsetY = param.filter.height/2 - Int(param.paddings[1]) - print("offset x: \(offsetX)") - print("offset y: \(offsetY)") + print(" param filter width: \(param.filter.width)") + print(" param filter height: \(param.filter.height)") + + print(" param paddings: \(param.paddings)") + + print("ConvBNReluKernel offset x: \(offsetX)") + print("ConvBNReluKernel offset y: \(offsetY)") let offsetZ = 0.0 @@ -116,8 +121,8 @@ class ConvBNReluKernel: Kernel, Computable, Testable { encoder.setTexture(param.output.metalTexture, index: 1) encoder.setBytes(&metalParam, length: MemoryLayout.size, index: 0) encoder.setBuffer(param.filter.buffer, offset: 0, index: 1) - encoder.setBuffer(param.newScale!, offset: 0, index: 3) - encoder.setBuffer(param.newBiase!, offset: 0, index: 4) + encoder.setBuffer(param.newScale!, offset: 0, index: 2) + encoder.setBuffer(param.newBiase!, offset: 0, index: 3) encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) encoder.endEncoding() } @@ -132,9 +137,8 @@ class ConvBNReluKernel: Kernel, Computable, Testable { var inMetalParam = param.metalParam encoder.setBytes(&inMetalParam, length: MemoryLayout.size, index: 0) encoder.setBuffer(param.filterBuffer, offset: 0, index: 1) - encoder.setBuffer(param.biaseBuffer, offset: 0, index: 2) - encoder.setBuffer(param.newScaleBuffer, offset: 0, index: 3) - encoder.setBuffer(param.newBiaseBuffer, offset: 0, index: 4) + encoder.setBuffer(param.newScaleBuffer, offset: 0, index: 2) + encoder.setBuffer(param.newBiaseBuffer, offset: 0, index: 3) encoder.dispatch(computePipline: pipline, outTexture: param.outputTexture) encoder.endEncoding() } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift index 029fdf0a45db067d874757bad38d91f1b6c70d59..09c80f9ab0e2e00becdc8794231f1341fc768df2 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift @@ -15,88 +15,99 @@ import Foundation struct PriorBoxMetalParam { - let offset: Float32 - let stepWidth: Float32 - let stepHeight: Float32 - let minSize: Float32 - let maxSize: Float32 - let imageWidth: Float32 - let imageHeight: Float32 - let clip: Bool - let numPriors: uint - let aspecRatiosSize: uint - let minSizeSize: uint - let maxSizeSize: uint + let offset: Float32 + let stepWidth: Float32 + let stepHeight: Float32 + let minSize: Float32 + let maxSize: Float32 + let imageWidth: Float32 + let imageHeight: Float32 + let clip: Bool + let numPriors: uint + let aspecRatiosSize: uint + let minSizeSize: uint + let maxSizeSize: uint } class PriorBoxKernel: Kernel, Computable{ - var metalParam: PriorBoxMetalParam! - - required init(device: MTLDevice, param: PriorBoxParam

) { - super.init(device: device, inFunctionName: "prior_box") - param.output.initTexture(device: device, transpose: [2, 0, 1, 3]) - param.outputVariances.initTexture(device: device, transpose: [2, 0, 1, 3]) - - let imageWidth = Float32(param.inputImage.originDim[3]) - let imageHeight = Float32(param.inputImage.originDim[2]) - - let featureWidth = param.inputImage.originDim[3] - let featureHeight = param.inputImage.originDim[2] - - if param.stepW == 0 || param.stepH == 0 { - param.stepW = Float32(imageWidth) / Float32(featureWidth) - param.stepH = Float32(imageHeight) / Float32(featureHeight) + var metalParam: PriorBoxMetalParam! + + required init(device: MTLDevice, param: PriorBoxParam

) { + super.init(device: device, inFunctionName: "prior_box") + param.output.initTexture(device: device, inTranspose: [2, 0, 1, 3]) + param.outputVariances.initTexture(device: device, inTranspose: [2, 0, 1, 3]) + + let imageWidth = Float32(param.inputImage.originDim[3]) + let imageHeight = Float32(param.inputImage.originDim[2]) + + let featureWidth = param.input.originDim[3] + let featureHeight = param.input.originDim[2] + + if param.stepW == 0 || param.stepH == 0 { + param.stepW = Float32(imageWidth) / Float32(featureWidth) + param.stepH = Float32(imageHeight) / Float32(featureHeight) + } + + var outputAspectRatior: [Float32] = [] + outputAspectRatior.append(1.0) + + let epsilon = 1e-6 + for ar in param.aspectRatios { + var alreadyExist = false + for outputAr in outputAspectRatior { + if fabs(Double(ar) - Double(outputAr)) < Double(epsilon) { + alreadyExist = true + break } + } - var outputAspectRatior: [Float32] = [] - outputAspectRatior.append(1.0) - - let epsilon = 1e-6 - for ar in param.aspectRatios { - var alreadyExist = false - for outputAr in outputAspectRatior { - if fabs(Double(ar) - Double(outputAr)) < Double(epsilon) { - alreadyExist = true - break - } - } - - if !alreadyExist { - outputAspectRatior.append(ar) - } - if param.flip { - outputAspectRatior.append(1.0 / ar) - } - } - - param.newAspectRatios = outputAspectRatior - let aspectRatiosSize = uint(outputAspectRatior.count) - - let maxSizeSize: uint = uint(param.maxSizes.count) - let minSizeSize: uint = uint(param.minSizes.count) - - let numPriors = aspectRatiosSize * minSizeSize + maxSizeSize - - let minSize = param.minSizes.last ?? 0.0 - let maxSize = param.maxSizes.last ?? 0.0 - - metalParam = PriorBoxMetalParam.init(offset: param.offset, stepWidth: param.stepW, stepHeight: param.stepH, minSize: minSize, maxSize: maxSize, imageWidth: imageWidth, imageHeight: imageHeight, clip: param.clip, numPriors: numPriors, aspecRatiosSize: aspectRatiosSize, minSizeSize: minSizeSize, maxSizeSize: maxSizeSize) - + if !alreadyExist { + outputAspectRatior.append(ar) + } + if param.flip { + outputAspectRatior.append(1.0 / ar) + } } - func compute(commandBuffer: MTLCommandBuffer, param: PriorBoxParam

) throws { - guard let encoder = commandBuffer.makeComputeCommandEncoder() else { - throw PaddleMobileError.predictError(message: " encode is nil") - } - encoder.setTexture(param.input.metalTexture, index: 0) - encoder.setTexture(param.output.metalTexture, index: 1) - encoder.setTexture(param.outputVariances.metalTexture, index: 2) - encoder.setBytes(&metalParam, length: MemoryLayout.size, index: 0) - encoder.setBytes(param.aspectRatios, length: MemoryLayout.size * param.aspectRatios.count, index: 1) - encoder.setBytes(param.variances, length: MemoryLayout.size * param.variances.count, index: 2) - encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) - encoder.endEncoding() + param.newAspectRatios = outputAspectRatior + let aspectRatiosSize = uint(outputAspectRatior.count) + + let maxSizeSize: uint = uint(param.maxSizes.count) + let minSizeSize: uint = uint(param.minSizes.count) + + let numPriors = aspectRatiosSize * minSizeSize + maxSizeSize + + let minSize = param.minSizes.last ?? 0.0 + let maxSize = param.maxSizes.last ?? 0.0 + + metalParam = PriorBoxMetalParam.init(offset: param.offset, stepWidth: param.stepW, stepHeight: param.stepH, minSize: minSize, maxSize: maxSize, imageWidth: imageWidth, imageHeight: imageHeight, clip: param.clip, numPriors: numPriors, aspecRatiosSize: aspectRatiosSize, minSizeSize: minSizeSize, maxSizeSize: maxSizeSize) + + } + + func compute(commandBuffer: MTLCommandBuffer, param: PriorBoxParam

) throws { + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { + throw PaddleMobileError.predictError(message: " encode is nil") } + print("metalParam: \(metalParam)") + + print(" newAspectRatios ") + print(param.newAspectRatios!) + + print(" clip: \(metalParam.clip)") + + print(" metalParam.numPriors: \(metalParam.numPriors)") + + print(" aspecRatiosSize: \(metalParam.aspecRatiosSize)") - + encoder.setTexture(param.input.metalTexture, index: 0) + encoder.setTexture(param.output.metalTexture, index: 1) + encoder.setTexture(param.outputVariances.metalTexture, index: 2) + encoder.setBytes(&metalParam, length: MemoryLayout.size, index: 0) + encoder.setBytes(param.newAspectRatios!, length: MemoryLayout.size * param.newAspectRatios!.count, index: 1) + encoder.setBytes(param.variances, length: MemoryLayout.size * param.variances.count, index: 2) + encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() + } + + } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Texture2DTo2DArrayKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Texture2DTo2DArrayKernel.swift index f82516459ea61bd61b35b953a61a4f75d1bd9629..8554beea2bc336294a3c577ccb4294e59f426bdd 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Texture2DTo2DArrayKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Texture2DTo2DArrayKernel.swift @@ -32,7 +32,7 @@ class Texture2DTo2DArrayKernel: Kernel, Computable{ } required init(device: MTLDevice, param: FeedParam

) { - param.output.initTexture(device: device, transpose: [0, 2, 3, 1]) + param.output.initTexture(device: device, inTranspose: [0, 2, 3, 1]) super.init(device: device, inFunctionName: "texture2d_to_2d_array") } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift index dd0577c878adffc4061044513e6ffcfa99e08908..33e1219b4d0fff972d8db3d16fc7ce1477841351 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift @@ -15,80 +15,92 @@ import Foundation struct TransposeMetalParam { - var iC: Int32 = 0 - var oC: Int32 = 0 - var i0: Int32 - var i1: Int32 - var i2: Int32 - var i3: Int32 - init(_ i0: Int32, _ i1: Int32, _ i2: Int32, _ i3: Int32) { - self.i0 = i0 - self.i1 = i1 - self.i2 = i2 - self.i3 = i3 - } - init(_ axis: [Int]) { - self.init(Int32(axis[0]), Int32(axis[1]), Int32(axis[2]), Int32(axis[3])) - } + var iC: Int32 = 0 + var oC: Int32 = 0 + var i0: Int32 + var i1: Int32 + var i2: Int32 + var i3: Int32 + init(_ i0: Int32, _ i1: Int32, _ i2: Int32, _ i3: Int32) { + self.i0 = i0 + self.i1 = i1 + self.i2 = i2 + self.i3 = i3 + } + init(_ axis: [Int]) { + self.init(Int32(axis[0]), Int32(axis[1]), Int32(axis[2]), Int32(axis[3])) + } } struct TransposeTestParam: TestParam { - let inputTexture: MTLTexture - let outputTexture: MTLTexture - let iC: Int - let oC: Int - let axis: [Int] + let inputTexture: MTLTexture + let outputTexture: MTLTexture + let iC: Int + let oC: Int + let axis: [Int] } class TransposeKernel: Kernel, Computable, Testable { - func compute(commandBuffer: MTLCommandBuffer, param: TransposeParam

) throws { - guard let encoder = commandBuffer.makeComputeCommandEncoder() else { - throw PaddleMobileError.predictError(message: " encode is nil") - } - var invT: [Int] = [0, 1, 2, 3] - for (i, v) in param.input.transpose.enumerated() { - invT[v] = i - } - var axis: [Int] = [0, 1, 2, 3] - for i in 0.. transpose! FAST :)") - } else { - print("====> transpose! SLOW :(") - } - encoder.setTexture(param.input.metalTexture, index: 0) - encoder.setTexture(param.output.metalTexture, index: 1) - encoder.setBytes(&tmp, length: MemoryLayout.size, index: 0) - encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) - encoder.endEncoding() + var metalParam: TransposeMetalParam! + func compute(commandBuffer: MTLCommandBuffer, param: TransposeParam

) throws { + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { + throw PaddleMobileError.predictError(message: " encode is nil") } + + encoder.setTexture(param.input.metalTexture, index: 0) + encoder.setTexture(param.output.metalTexture, index: 1) + encoder.setBytes(&metalParam, length: MemoryLayout.size, index: 0) + encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() + } + + required init(device: MTLDevice, param: TransposeParam

) { + param.output.initTexture(device: device, inTranspose: [0, 1, 2, 3]) + super.init(device: device, inFunctionName: "transpose") - required init(device: MTLDevice, param: TransposeParam

) { - param.output.initTexture(device: device, transpose: [0, 1, 2, 3]) - super.init(device: device, inFunctionName: "transpose") + var invT: [Int] = [0, 1, 2, 3] + for (i, v) in param.input.transpose.enumerated() { + invT[v] = i } - required init(device: MTLDevice, testParam: TransposeTestParam) { - super.init(device: device, inFunctionName: "transpose") + var axis: [Int] = [0, 1, 2, 3] + +// var doNothing = false +// if param.axis.count == param.input.transpose.count { +// doNothing = param.axis == param.input.transpose.map { Int32($0) } +// } + + + for i in 0.. transpose! FAST :)") + } else { + print("====> transpose! SLOW :(") + } + metalParam = tmp + } + required init(device: MTLDevice, testParam: TransposeTestParam) { + super.init(device: device, inFunctionName: "transpose") + fatalError() + } + + public func test(commandBuffer: MTLCommandBuffer, param: TransposeTestParam) { + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { + fatalError() } - public func test(commandBuffer: MTLCommandBuffer, param: TransposeTestParam) { - guard let encoder = commandBuffer.makeComputeCommandEncoder() else { - fatalError() - } - - encoder.setTexture(param.inputTexture, index: 0) - encoder.setTexture(param.outputTexture, index: 1) - var tmp = TransposeMetalParam.init(param.axis) - tmp.iC = Int32(param.iC) - tmp.oC = Int32(param.oC) - - encoder.setBytes(&tmp, length: MemoryLayout.size, index: 0) - encoder.dispatch(computePipline: pipline, outTexture: param.outputTexture) - encoder.endEncoding() - }} + encoder.setTexture(param.inputTexture, index: 0) + encoder.setTexture(param.outputTexture, index: 1) + var tmp = TransposeMetalParam.init(param.axis) + tmp.iC = Int32(param.iC) + tmp.oC = Int32(param.oC) + + encoder.setBytes(&tmp, length: MemoryLayout.size, index: 0) + encoder.dispatch(computePipline: pipline, outTexture: param.outputTexture) + encoder.endEncoding() + }} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Common.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Common.metal new file mode 100644 index 0000000000000000000000000000000000000000..a4e9de2bd4f12b67dc947776bf4f7e9ab01618b1 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Common.metal @@ -0,0 +1,55 @@ +// +// common.metal +// paddle-mobile +// +// Created by liuRuiLong on 2018/8/26. +// Copyright © 2018年 orange. All rights reserved. +// + +#include +using namespace metal; + + + +inline void xyzn2abcd(int C, int xyzn[4], int abcd[4]) { + abcd[2] = xyzn[0]; + abcd[1] = xyzn[1]; + uint t = xyzn[2] * 4 + xyzn[3]; + abcd[0] = t / C; + abcd[3] = t % C; +} + +inline void abcd2xyzn(int C, int abcd[4], int xyzn[4]) { + xyzn[0] = abcd[2]; + xyzn[1] = abcd[1]; + uint t = abcd[0] * C + abcd[3]; + xyzn[2] = t / 4; + xyzn[3] = t % 4; +} + +inline int32_t abcd2index(int32_t dim[4], int32_t abcd[4]) { + int32_t r = abcd[0]; + r = r * dim[1] + abcd[1]; + r = r * dim[2] + abcd[2]; + r = r * dim[3] + abcd[3]; + return r; +} + +inline void index2abcd(int32_t dim[4], int32_t ind, int32_t abcd[4]) { + abcd[3] = ind % dim[3]; ind /= dim[3]; + abcd[2] = ind % dim[2]; ind /= dim[2]; + abcd[1] = ind % dim[1]; ind /= dim[1]; + abcd[0] = ind; +} + +inline void trans(int32_t trans[4], int32_t ipos[4], int32_t opos[4]) { + for (int i = 0; i < 4; i++) { + opos[i] = ipos[trans[i]]; + } +} + +inline void invtrans(int32_t trans[4], int32_t ipos[4], int32_t opos[4]) { + for (int i = 0; i < 4; i++) { + opos[trans[i]] = ipos[i]; + } +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ConvKernel.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ConvKernel.metal index 5a059d89bcc98d9be658d8f3346f31ad747b68d9..15f6704a90a7e15e8140ae1f6fb4ddc0eae9286a 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ConvKernel.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ConvKernel.metal @@ -704,9 +704,8 @@ kernel void conv_batch_norm_relu_1x1(texture2d_array inTe texture2d_array outTexture [[texture(1)]], constant MetalConvParam ¶m [[buffer(0)]], const device float4 *weights [[buffer(1)]], - const device float4 *biase [[buffer(2)]], - const device float4 *new_scale [[buffer(3)]], - const device float4 *new_biase [[buffer(4)]], + const device float4 *new_scale [[buffer(2)]], + const device float4 *new_biase [[buffer(3)]], uint3 gid [[thread_position_in_grid]]) { if (gid.x >= outTexture.get_width() || @@ -749,9 +748,8 @@ kernel void conv_batch_norm_relu_3x3(texture2d_array inTe texture2d_array outTexture [[texture(1)]], constant MetalConvParam ¶m [[buffer(0)]], const device float4 *weights [[buffer(1)]], - const device float4 *biase [[buffer(2)]], - const device float4 *new_scale [[buffer(3)]], - const device float4 *new_biase [[buffer(4)]], + const device float4 *new_scale [[buffer(2)]], + const device float4 *new_biase [[buffer(3)]], uint3 gid [[thread_position_in_grid]]) { if (gid.x >= outTexture.get_width() || @@ -803,8 +801,8 @@ kernel void depthwise_conv_batch_norm_relu_3x3(texture2d_array outTexture [[texture(1)]], constant MetalConvParam ¶m [[buffer(0)]], const device float *weights [[buffer(1)]], - const device float4 *new_scale [[buffer(3)]], - const device float4 *new_biase [[buffer(4)]], + const device float4 *new_scale [[buffer(2)]], + const device float4 *new_biase [[buffer(3)]], uint3 gid [[thread_position_in_grid]]) { if (gid.x >= outTexture.get_width() || diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Kernels.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Kernels.metal index 8c2566d71d9c654b1ddf55ff2769994c370949d8..4032d0c364d4422d8f6c1bcc5f4a8cf0f6c486b8 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Kernels.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/Kernels.metal @@ -13,39 +13,40 @@ limitations under the License. */ #include +#include "Common.metal" using namespace metal; struct OutputDim { - ushort width; - ushort height; - ushort strideX; - ushort strideY; + ushort width; + ushort height; + ushort strideX; + ushort strideY; }; kernel void resize(texture2d inTexture [[texture(0)]], texture2d_array outTexture [[texture(1)]], constant OutputDim ¶ms [[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; - - constexpr sampler s(coord::pixel, filter::nearest, address::clamp_to_zero); - const uint2 pos = gid.xy * uint2(params.strideX, params.strideY); - const half4 input = inTexture.read(pos); - outTexture.write(half4(input.x, input.y, input.z, input.w), gid.xy, gid.z); + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height() || + gid.z >= outTexture.get_array_size()) return; + + constexpr sampler s(coord::pixel, filter::nearest, address::clamp_to_zero); + const uint2 pos = gid.xy * uint2(params.strideX, params.strideY); + const half4 input = inTexture.read(pos); + outTexture.write(half4(input.x, input.y, input.z, input.w), gid.xy, gid.z); } kernel void elementwise_add(texture2d_array inTexture [[texture(0)]], texture2d_array outTexture [[texture(1)]], const device half4 *biasTerms [[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; - constexpr sampler s(coord::pixel, filter::nearest, address::clamp_to_zero); - const half4 input = inTexture.read(gid.xy, gid.z); - outTexture.write(input, gid.xy, gid.z); + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height() || + gid.z >= outTexture.get_array_size()) return; + constexpr sampler s(coord::pixel, filter::nearest, address::clamp_to_zero); + const half4 input = inTexture.read(gid.xy, gid.z); + outTexture.write(input, gid.xy, gid.z); } @@ -63,280 +64,190 @@ kernel void elementwise_add(texture2d_array inTexture [[text kernel void texture2d_to_2d_array(texture2d inTexture [[texture(0)]], texture2d_array outTexture [[texture(1)]], uint3 gid [[thread_position_in_grid]]) { - if (gid.x >= inTexture.get_width() || - gid.y >= inTexture.get_height()){ - return; - } - const float4 input = inTexture.read(gid.xy); - outTexture.write(input, gid.xy, 0); + if (gid.x >= inTexture.get_width() || + gid.y >= inTexture.get_height()){ + return; + } + const float4 input = inTexture.read(gid.xy); + outTexture.write(input, gid.xy, 0); } kernel void texture2d_to_2d_array_half(texture2d inTexture [[texture(0)]], - texture2d_array outTexture [[texture(1)]], - uint3 gid [[thread_position_in_grid]]) { - if (gid.x >= inTexture.get_width() || - gid.y >= inTexture.get_height()){ - return; - } - const half4 input = inTexture.read(gid.xy); - outTexture.write(input, gid.xy, 0); + texture2d_array outTexture [[texture(1)]], + uint3 gid [[thread_position_in_grid]]) { + if (gid.x >= inTexture.get_width() || + gid.y >= inTexture.get_height()){ + return; + } + const half4 input = inTexture.read(gid.xy); + outTexture.write(input, gid.xy, 0); } struct PoolParam { - int ksizeX; - int ksizeY; - int strideX; - int strideY; - int paddingX; - int paddingY; - int poolType; + int ksizeX; + int ksizeY; + int strideX; + int strideY; + int paddingX; + int paddingY; + 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; + 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)); + } } - outTexture.write(r, gid.xy, 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); } 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; + 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)); + } } - outTexture.write(r, gid.xy, 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); } kernel void softmax(texture2d_array inTexture [[texture(0)]], texture2d_array outTexture [[texture(1)]], 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 zsize = inTexture.get_array_size(); - float maxv = inTexture.read(uint2(0, 0), 0)[0]; - for (int z = 0; z < zsize; z++) { - float4 r = inTexture.read(uint2(0, 0), z); - maxv = max(maxv, max(max(r[0], r[1]), max(r[2], r[3]))); - } - float sum = 0; - for (int z = 0; z < zsize; z++) { - float4 r = inTexture.read(uint2(0, 0), z); - sum += exp(r[0] - maxv) + exp(r[1] - maxv) + exp(r[2] - maxv) + exp(r[3] - maxv); - } - float4 rr = inTexture.read(gid.xy, gid.z); - rr = exp(rr - maxv) / sum; - outTexture.write(rr, gid.xy, gid.z); + if (gid.x >= outTexture.get_width() || + gid.y >= outTexture.get_height() || + gid.z >= outTexture.get_array_size()) return; + int zsize = inTexture.get_array_size(); + float maxv = inTexture.read(uint2(0, 0), 0)[0]; + for (int z = 0; z < zsize; z++) { + float4 r = inTexture.read(uint2(0, 0), z); + maxv = max(maxv, max(max(r[0], r[1]), max(r[2], r[3]))); + } + float sum = 0; + for (int z = 0; z < zsize; z++) { + float4 r = inTexture.read(uint2(0, 0), z); + sum += exp(r[0] - maxv) + exp(r[1] - maxv) + exp(r[2] - maxv) + exp(r[3] - maxv); + } + float4 rr = inTexture.read(gid.xy, gid.z); + rr = exp(rr - maxv) / sum; + outTexture.write(rr, gid.xy, gid.z); } kernel void softmax_half(texture2d_array inTexture [[texture(0)]], - texture2d_array outTexture [[texture(1)]], - 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 zsize = inTexture.get_array_size(); - half maxv = inTexture.read(uint2(0, 0), 0)[0]; - for (int z = 0; z < zsize; z++) { - half4 r = inTexture.read(uint2(0, 0), z); - maxv = max(maxv, max(max(r[0], r[1]), max(r[2], r[3]))); - } - float sum = 0; - for (int z = 0; z < zsize; z++) { - half4 r = inTexture.read(uint2(0, 0), z); - sum += exp(r[0] - maxv) + exp(r[1] - maxv) + exp(r[2] - maxv) + exp(r[3] - maxv); - } - half4 rr = inTexture.read(gid.xy, gid.z); - rr = exp(rr - maxv) / sum; - outTexture.write(rr, gid.xy, gid.z); + texture2d_array outTexture [[texture(1)]], + 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 zsize = inTexture.get_array_size(); + half maxv = inTexture.read(uint2(0, 0), 0)[0]; + for (int z = 0; z < zsize; z++) { + half4 r = inTexture.read(uint2(0, 0), z); + maxv = max(maxv, max(max(r[0], r[1]), max(r[2], r[3]))); + } + float sum = 0; + for (int z = 0; z < zsize; z++) { + half4 r = inTexture.read(uint2(0, 0), z); + sum += exp(r[0] - maxv) + exp(r[1] - maxv) + exp(r[2] - maxv) + exp(r[3] - maxv); + } + half4 rr = inTexture.read(gid.xy, gid.z); + rr = exp(rr - maxv) / sum; + outTexture.write(rr, gid.xy, gid.z); } -inline void xyzn2abcd(int C, int xyzn[4], int abcd[4]) { - abcd[2] = xyzn[0]; - abcd[1] = xyzn[1]; - uint t = xyzn[2] * 4 + xyzn[3]; - abcd[0] = t / C; - abcd[3] = t % C; -} -inline void abcd2xyzn(int C, int abcd[4], int xyzn[4]) { - xyzn[0] = abcd[2]; - xyzn[1] = abcd[1]; - uint t = abcd[0] * C + abcd[3]; - xyzn[2] = t / 4; - xyzn[3] = t % 4; -} - -inline int32_t abcd2index(int32_t dim[4], int32_t abcd[4]) { - int32_t r = abcd[0]; - r = r * dim[1] + abcd[1]; - r = r * dim[2] + abcd[2]; - r = r * dim[3] + abcd[3]; - return r; -} - -inline void index2abcd(int32_t dim[4], int32_t ind, int32_t abcd[4]) { - abcd[3] = ind % dim[3]; ind /= dim[3]; - abcd[2] = ind % dim[2]; ind /= dim[2]; - abcd[1] = ind % dim[1]; ind /= dim[1]; - abcd[0] = ind; -} - -inline void trans(int32_t trans[4], int32_t ipos[4], int32_t opos[4]) { - for (int i = 0; i < 4; i++) { - opos[i] = ipos[trans[i]]; - } -} - -inline void invtrans(int32_t trans[4], int32_t ipos[4], int32_t opos[4]) { - for (int i = 0; i < 4; i++) { - opos[trans[i]] = ipos[i]; - } -} struct TransposeParam { - int iC; - int oC; - int axis[4]; + int iC; + int oC; + int axis[4]; }; kernel void transpose(texture2d_array inTexture [[texture(0)]], - texture2d_array outTexture [[texture(1)]], - constant TransposeParam &pm [[buffer(0)]], - uint3 gid [[thread_position_in_grid]]) { - - if ((pm.axis[0] == 0) && (pm.axis[1] == 1) && (pm.axis[2] == 2) && (pm.axis[3] == 3)) { - // do nothing - float4 r = inTexture.read(gid.xy, gid.z); - outTexture.write(r, gid.xy, gid.z); - } else { - float4 r; - for (int n = 0; n < 4; n++) { - int ixyzn[] = {int(gid.x), int(gid.y), int(gid.z), n}; - int iabcd[4], oabcd[4], oxyzn[4]; - xyzn2abcd(pm.oC, ixyzn, iabcd); - oabcd[pm.axis[0]] = iabcd[0]; - oabcd[pm.axis[1]] = iabcd[1]; - oabcd[pm.axis[2]] = iabcd[2]; - oabcd[pm.axis[3]] = iabcd[3]; - abcd2xyzn(pm.iC, oabcd, oxyzn); - float4 rt = inTexture.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2]); - r[n] = rt[oxyzn[3]]; - } - outTexture.write(r, gid.xy, gid.z); - } -} - -struct ReshapeParam { - int32_t idim[4]; - int32_t itrans[4]; - int32_t odim[4]; - int32_t otrans[4]; -}; + texture2d_array outTexture [[texture(1)]], + constant TransposeParam &pm [[buffer(0)]], + uint3 gid [[thread_position_in_grid]]) { + -kernel void reshape(texture2d_array inTexture [[texture(0)]], - texture2d_array outTexture [[texture(1)]], - constant ReshapeParam &rp [[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 oxyzn[4] = {int(gid.x), int(gid.y), int(gid.z), 0}, oabcd[4], ixyzn[4], iabcd[4]; - ReshapeParam lrp = rp; - int oC = lrp.odim[lrp.otrans[3]]; - int iC = lrp.idim[lrp.itrans[3]]; - int count = lrp.odim[0] * lrp.odim[1] * lrp.odim[2] * lrp.odim[3]; + if ((pm.axis[0] == 0) && (pm.axis[1] == 1) && (pm.axis[2] == 2) && (pm.axis[3] == 3)) { + // do nothing + float4 r = inTexture.read(gid.xy, gid.z); + outTexture.write(r, gid.xy, gid.z); + } else { float4 r; for (int n = 0; n < 4; n++) { - oxyzn[3] = n; - xyzn2abcd(oC, oxyzn, oabcd); - int tabcd[4]; - invtrans(lrp.otrans, oabcd, tabcd); - int index = abcd2index(lrp.odim, tabcd); - if (index < count) { - index2abcd(lrp.idim, index, tabcd); - trans(lrp.itrans, tabcd, iabcd); - abcd2xyzn(iC, tabcd, ixyzn); - r[n] = inTexture.read(uint2(ixyzn[0], ixyzn[1]), ixyzn[2])[ixyzn[3]]; - } else { - r[n] = 0; - } + int ixyzn[] = {int(gid.x), int(gid.y), int(gid.z), n}; + int iabcd[4], oabcd[4], oxyzn[4]; + xyzn2abcd(pm.oC, ixyzn, iabcd); + oabcd[pm.axis[0]] = iabcd[0]; + oabcd[pm.axis[1]] = iabcd[1]; + oabcd[pm.axis[2]] = iabcd[2]; + oabcd[pm.axis[3]] = iabcd[3]; + abcd2xyzn(pm.iC, oabcd, oxyzn); + float4 rt = inTexture.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2]); + r[n] = rt[oxyzn[3]]; } outTexture.write(r, gid.xy, gid.z); + } } -// -//kernel void reshape_half(texture2d_array inTexture [[texture(0)]], -// texture2d_array outTexture [[texture(1)]], -// uint3 gid [[thread_position_in_grid]]) { -// if (gid.x >= outTexture.get_width() || -// gid.y >= outTexture.get_height() || -// gid.z >= outTexture.get_array_size()) return; -// -// half4 r = inTexture.read(uint2(0, 0), gid.x); -// outTexture.write(r, gid.xy, gid.z); -//} struct ConcatParam { - int32_t odim[4]; - int32_t axis; - int32_t offset; - int32_t vdim[6]; + int32_t odim[4]; + int32_t axis; + int32_t offset; + int32_t vdim[6]; }; kernel void concat(texture2d_array in0 [[texture(0)]], @@ -349,39 +260,39 @@ kernel void concat(texture2d_array in0 [[texture(0)]], texture2d_array out [[texture(7)]], constant ConcatParam & pm [[buffer(0)]], uint3 gid [[thread_position_in_grid]]) { - ConcatParam cp = pm; - int xyzn[4] = {int(gid.x), int(gid.y), int(gid.z), 0}, abcd[4], oxyzn[4]; - float4 r; - for (int i = 0; i < 4; i++) { - xyzn[3] = i; - xyzn2abcd(cp.odim[3], xyzn, abcd); - int k = abcd[cp.axis] - cp.offset; - int j = 0; - if (k < 0) { - r[i] = inx.read(gid.xy, gid.z)[i]; - } else { - for (; j < 6; j++) { - if (k < cp.vdim[j]) { - break; - } - k -= cp.vdim[j]; - } - int ta = cp.odim[cp.axis]; - abcd[cp.axis] = k; - cp.odim[cp.axis] = cp.vdim[j]; - abcd2xyzn(cp.odim[3], abcd, oxyzn); - cp.odim[cp.axis] = ta; - switch (j) { - case 0: r[i] = in0.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; - case 1: r[i] = in1.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; - case 2: r[i] = in2.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; - case 3: r[i] = in3.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; - case 4: r[i] = in4.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; - case 5: r[i] = in5.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; - } + ConcatParam cp = pm; + int xyzn[4] = {int(gid.x), int(gid.y), int(gid.z), 0}, abcd[4], oxyzn[4]; + float4 r; + for (int i = 0; i < 4; i++) { + xyzn[3] = i; + xyzn2abcd(cp.odim[3], xyzn, abcd); + int k = abcd[cp.axis] - cp.offset; + int j = 0; + if (k < 0) { + r[i] = inx.read(gid.xy, gid.z)[i]; + } else { + for (; j < 6; j++) { + if (k < cp.vdim[j]) { + break; } + k -= cp.vdim[j]; + } + int ta = cp.odim[cp.axis]; + abcd[cp.axis] = k; + cp.odim[cp.axis] = cp.vdim[j]; + abcd2xyzn(cp.odim[3], abcd, oxyzn); + cp.odim[cp.axis] = ta; + switch (j) { + case 0: r[i] = in0.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; + case 1: r[i] = in1.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; + case 2: r[i] = in2.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; + case 3: r[i] = in3.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; + case 4: r[i] = in4.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; + case 5: r[i] = in5.read(uint2(oxyzn[0], oxyzn[1]), oxyzn[2])[oxyzn[3]]; break; + } } - out.write(r, gid.xy, gid.z); + } + out.write(r, gid.xy, gid.z); } kernel void boxcoder(texture2d_array priorBox [[texture(0)]], @@ -389,12 +300,12 @@ kernel void boxcoder(texture2d_array priorBox [[texture(0)] texture2d_array targetBox [[texture(2)]], texture2d_array output[[texture(3)]], uint3 gid [[thread_position_in_grid]]) { - float4 t = targetBox.read(gid.xy, gid.z); - float4 p = priorBox.read(gid.xy, gid.z); - float4 pv = priorBoxVar.read(gid.xy, gid.z); - float ox = (p.z * pv.x * t.x + p.x) - t.z / 2; - float oy = (p.w * pv.y * t.y + p.y) - t.w / 2; - float ow = exp(pv.z * t.z) * p.z + t.z / 2; - float oh = exp(pv.w * t.w) * p.w + t.w / 2; - output.write(float4(ox, oy, ow, oh), gid.xy, gid.z); + float4 t = targetBox.read(gid.xy, gid.z); + float4 p = priorBox.read(gid.xy, gid.z); + float4 pv = priorBoxVar.read(gid.xy, gid.z); + float ox = (p.z * pv.x * t.x + p.x) - t.z / 2; + float oy = (p.w * pv.y * t.y + p.y) - t.w / 2; + float ow = exp(pv.z * t.z) * p.z + t.z / 2; + float oh = exp(pv.w * t.w) * p.w + t.w / 2; + output.write(float4(ox, oy, ow, oh), gid.xy, gid.z); } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PriorBoxKernel.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PriorBoxKernel.metal index 7c6ab6dd03a512b99f391f7afebe3a8f7999a9d5..6083c6b514a3d8a0918d585a950d915e69a045fe 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PriorBoxKernel.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/PriorBoxKernel.metal @@ -60,7 +60,7 @@ kernel void prior_box(texture2d_array inTexture [[texture(0 float4 res; if (param.clip) { - res = min(max(box, 0.0), 1.0); + res = fmin(fmax(box, 0.0), 1.0); } else { res = box; } @@ -74,7 +74,7 @@ kernel void prior_box(texture2d_array inTexture [[texture(0 max_box.y = (center_y - box_height) / param.imageHeight; max_box.z = (center_x + box_width) / param.imageWidth; max_box.w = (center_y + box_height) / param.imageHeight; - + float4 res; if (param.clip) { res = min(max(max_box, 0.0), 1.0); @@ -92,6 +92,7 @@ kernel void prior_box(texture2d_array inTexture [[texture(0 variances_output.y = variance.y; variances_output.z = variance.z; variances_output.w = variance.w; + varianceTexture.write(variances_output, gid.xy, gid.z); } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal new file mode 100644 index 0000000000000000000000000000000000000000..e9ea49f1d774f77a9eabff7f6b1095329395d01a --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal @@ -0,0 +1,82 @@ +/* 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 +#include "Common.metal" + +using namespace metal; + +struct ReshapeParam { + int32_t idim[4]; + int32_t itrans[4]; + int32_t odim[4]; + int32_t otrans[4]; +}; + +kernel void reshape(texture2d_array inTexture [[texture(0)]], + texture2d_array outTexture [[texture(1)]], + constant ReshapeParam &rp [[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 oxyzn[4] = {int(gid.x), int(gid.y), int(gid.z), 0}, oabcd[4], ixyzn[4]; + ReshapeParam lrp = rp; + int oC = lrp.odim[lrp.otrans[3]]; + int iC = lrp.idim[lrp.itrans[3]]; + int count = lrp.odim[0] * lrp.odim[1] * lrp.odim[2] * lrp.odim[3]; + float4 r; + for (int n = 0; n < 4; n++) { + oxyzn[3] = n; + + //4 (gid.x gid.y, gid.z, 0~4) + xyzn2abcd(oC, oxyzn, oabcd); + int tabcd[4]; + invtrans(lrp.otrans, oabcd, tabcd); + int index = abcd2index(lrp.odim, tabcd); + if (index < count) { + int c = index % 4; + + int temp0 = index % (inTexture.get_array_size() * 4); + int slice = temp0 / 4; + + int temp1 = index % (inTexture.get_array_size() * 4 * lrp.idim[2]); + int w = temp1 / (inTexture.get_array_size() * 4); + + int h = index / (inTexture.get_array_size() * 4 * lrp.idim[2]); + + +// index2abcd(lrp.idim, index, tabcd); +// abcd2xyzn(iC, tabcd, ixyzn); + r[n] = inTexture.read(uint2(w, h), slice)[c]; + } else { + r[n] = 0; + } + } + outTexture.write(r, gid.xy, gid.z); +} +// +//kernel void reshape_half(texture2d_array inTexture [[texture(0)]], +// texture2d_array outTexture [[texture(1)]], +// uint3 gid [[thread_position_in_grid]]) { +// if (gid.x >= outTexture.get_width() || +// gid.y >= outTexture.get_height() || +// gid.z >= outTexture.get_array_size()) return; +// +// half4 r = inTexture.read(uint2(0, 0), gid.x); +// outTexture.write(r, gid.xy, gid.z); +//} + + diff --git a/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift index 7931d3b5611546e3e9f06fa556c061fa85fe6804..590427d9eff883a069f8a96509db50b01f4272e9 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift @@ -31,7 +31,11 @@ class MulticlassNMSParam: OpParam { } class MulticlassNMSOp: Operator, MulticlassNMSParam

>, Runable, Creator, InferShaperable{ - + + func inputs() -> [Variant] { + return [para.scores,para.bboxes] + } + func inferShape() { // para.output.dim = para.input.dim } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift index 07676defe71ec18560df4be630cd04008cd1aad6..0c2ae767f80a96ca0ea7d03014050d23a1987cb6 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/PoolOp.swift @@ -15,54 +15,58 @@ import Foundation class PoolParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - input = try PoolParam.inputX(inputs: opDesc.inputs, from: inScope) - output = try PoolParam.outputOut(outputs: opDesc.outputs, from: inScope) - poolType = try PoolParam.getAttr(key: "pooling_type", attrs: opDesc.attrs) - ksize = try PoolParam.getAttr(key: "ksize", attrs: opDesc.attrs) - stride = try PoolParam.getAttr(key: "strides", attrs: opDesc.attrs) - padding = try PoolParam.getAttr(key: "paddings", attrs: opDesc.attrs) - ceilMode = try PoolParam.getAttr(key: "ceil_mode", attrs: opDesc.attrs) - globalPooling = try PoolParam.getAttr(key: "global_pooling", attrs: opDesc.attrs) - } catch let error { - throw error - } -// let buffer = input.metalTexture.buffer.contents().assumingMemoryBound(to: P.self) + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + input = try PoolParam.inputX(inputs: opDesc.inputs, from: inScope) + output = try PoolParam.outputOut(outputs: opDesc.outputs, from: inScope) + poolType = try PoolParam.getAttr(key: "pooling_type", attrs: opDesc.attrs) + ksize = try PoolParam.getAttr(key: "ksize", attrs: opDesc.attrs) + stride = try PoolParam.getAttr(key: "strides", attrs: opDesc.attrs) + padding = try PoolParam.getAttr(key: "paddings", attrs: opDesc.attrs) + ceilMode = try PoolParam.getAttr(key: "ceil_mode", attrs: opDesc.attrs) + globalPooling = try PoolParam.getAttr(key: "global_pooling", attrs: opDesc.attrs) + } catch let error { + throw error } - let input: Texture

- var output: Texture

- var ksize: [Int32] - var stride: [Int32] - var padding: [Int32] - var poolType: String - var ceilMode: Bool - var globalPooling: Bool + // let buffer = input.metalTexture.buffer.contents().assumingMemoryBound(to: P.self) + } + let input: Texture

+ var output: Texture

+ var ksize: [Int32] + var stride: [Int32] + var padding: [Int32] + var poolType: String + var ceilMode: Bool + var globalPooling: Bool } class PoolOp: Operator, PoolParam

>, Runable, Creator, InferShaperable{ - - func inferShape() { - // para.output.dim = para.input.dim - } - - typealias OpType = PoolOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } - } - - func delogOutput() { - print("pool2d delog") - 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: true) + + func inputs() -> [Variant] { + return [para.input] + } + + func inferShape() { + // para.output.dim = para.input.dim + } + + typealias OpType = PoolOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } + } + + func delogOutput() { + print("pool2d delog") + 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: true) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/PreluOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/PreluOp.swift index 26d19db2867037441832e028a578cca24769eeb2..1b55c1f3cbbb417ebd36d914b563ae3c491f3565 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/PreluOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/PreluOp.swift @@ -35,6 +35,10 @@ class PreluParam: OpParam { class PreluOp: Operator, PreluParam

>, Runable, Creator, InferShaperable{ + func inputs() -> [Variant] { + return [para.alpha, para.input] + } + func inferShape() { // para.output.dim = para.input.dim } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift index 7b6817ee06e04947ce53f3dbdbc33ca6e54dedd0..8610ab7a90bd444f9c79c1c16c53dad25d9efb00 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift @@ -15,57 +15,88 @@ import Foundation class PriorBoxParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - input = try PriorBoxParam.input(inputs: opDesc.inputs, from: inScope) - output = try PriorBoxParam.outputBoxes(outputs: opDesc.outputs, from: inScope) - inputImage = try PriorBoxParam.inputImage(inputs: opDesc.inputs, from: inScope) - outputVariances = try PriorBoxParam.outputVariances(outputs: opDesc.outputs, from: inScope) - minSizes = try PriorBoxParam.getAttr(key: "min_sizes", attrs: opDesc.attrs) - maxSizes = try PriorBoxParam.getAttr(key: "max_sizes", attrs: opDesc.attrs) - aspectRatios = try PriorBoxParam.getAttr(key: "aspect_ratios", attrs: opDesc.attrs) - variances = try PriorBoxParam.getAttr(key: "variances", attrs: opDesc.attrs) - flip = try PriorBoxParam.getAttr(key: "flip", attrs: opDesc.attrs) - clip = try PriorBoxParam.getAttr(key: "clip", attrs: opDesc.attrs) - stepW = try PriorBoxParam.getAttr(key: "step_w", attrs: opDesc.attrs) - stepH = try PriorBoxParam.getAttr(key: "step_h", attrs: opDesc.attrs) - offset = try PriorBoxParam.getAttr(key: "offset", attrs: opDesc.attrs) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + input = try PriorBoxParam.input(inputs: opDesc.inputs, from: inScope) + output = try PriorBoxParam.outputBoxes(outputs: opDesc.outputs, from: inScope) + inputImage = try PriorBoxParam.inputImage(inputs: opDesc.inputs, from: inScope) + outputVariances = try PriorBoxParam.outputVariances(outputs: opDesc.outputs, from: inScope) + minSizes = try PriorBoxParam.getAttr(key: "min_sizes", attrs: opDesc.attrs) + maxSizes = try PriorBoxParam.getAttr(key: "max_sizes", attrs: opDesc.attrs) + aspectRatios = try PriorBoxParam.getAttr(key: "aspect_ratios", attrs: opDesc.attrs) + variances = try PriorBoxParam.getAttr(key: "variances", attrs: opDesc.attrs) + flip = try PriorBoxParam.getAttr(key: "flip", attrs: opDesc.attrs) + clip = try PriorBoxParam.getAttr(key: "clip", attrs: opDesc.attrs) + stepW = try PriorBoxParam.getAttr(key: "step_w", attrs: opDesc.attrs) + stepH = try PriorBoxParam.getAttr(key: "step_h", attrs: opDesc.attrs) + offset = try PriorBoxParam.getAttr(key: "offset", attrs: opDesc.attrs) + } catch let error { + throw error } - - let minSizes: [Float32] - let maxSizes: [Float32] - let aspectRatios: [Float32] - var newAspectRatios: [Float32]? - let variances: [Float32] - let flip: Bool - let clip: Bool - var stepW: Float32 - var stepH: Float32 - let offset: Float32 - - let input: Texture

- let inputImage: Texture

- var output: Texture

- let outputVariances: Texture

+ } + + let minSizes: [Float32] + let maxSizes: [Float32] + let aspectRatios: [Float32] + var newAspectRatios: [Float32]? + let variances: [Float32] + let flip: Bool + let clip: Bool + var stepW: Float32 + var stepH: Float32 + let offset: Float32 + + let input: Texture

+ let inputImage: Texture

+ var output: Texture

+ let outputVariances: Texture

} class PriorBoxOp: Operator, PriorBoxParam

>, Runable, Creator, InferShaperable{ - - func inferShape() { + + func inputs() -> [Variant] { + return [para.input, para.inputImage] + } + + + func inferShape() { + } + + typealias OpType = PriorBoxOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } - - typealias OpType = PriorBoxOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } + } + + func delogOutput() { + print("pribox: ") + print("output: ") + // output + let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in + return o } + + print(outputArray) + +// writeToLibrary(fileName: "box_out", array: outputArray) + + // output variance +// let outputVarianceArray = para.outputVariances.metalTexture.floatArray { (o: Float32) -> Float32 in +// return o +// } +// +// print(" output variance: \(outputVarianceArray)") + +// writeToLibrary(fileName: "variance_out", array: outputVarianceArray) + + print("pribox write done ") + + + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ReluOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ReluOp.swift index f65e402cdd2b6356199a2104f99556cd4fdd3b6a..5d75f8ae816a9af6da1f5e2b480667b30116aa06 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ReluOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ReluOp.swift @@ -15,33 +15,37 @@ import Foundation class ReluParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - input = try ReluParam.inputX(inputs: opDesc.inputs, from: inScope) - output = try ReluParam.outputOut(outputs: opDesc.outputs, from: inScope) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + input = try ReluParam.inputX(inputs: opDesc.inputs, from: inScope) + output = try ReluParam.outputOut(outputs: opDesc.outputs, from: inScope) + } catch let error { + throw error } - let input: Texture

- var output: Texture

+ } + let input: Texture

+ var output: Texture

} class ReluOp: Operator, ReluParam

>, Runable, Creator, InferShaperable{ - - func inferShape() { - para.output.dim = para.input.dim - } - - typealias OpType = ReluOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } + + func inputs() -> [Variant] { + return [para.input] + } + + func inferShape() { + para.output.dim = para.input.dim + } + + typealias OpType = ReluOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift index 5ef52c407c9ed8268f7117a6b52ef24211eb6bb3..b249dc3883ab9110fe3844eeac4a49ba2f868ef1 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift @@ -15,40 +15,44 @@ import Foundation class ReshapeParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - input = try ReshapeParam.inputX(inputs: opDesc.inputs, from: inScope) - output = try ReshapeParam.outputOut(outputs: opDesc.outputs, from: inScope) -// shape = output.dim - inplace = try ReshapeParam.getAttr(key: "inplace", attrs: opDesc.attrs) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + input = try ReshapeParam.inputX(inputs: opDesc.inputs, from: inScope) + output = try ReshapeParam.outputOut(outputs: opDesc.outputs, from: inScope) + // shape = output.dim + inplace = try ReshapeParam.getAttr(key: "inplace", attrs: opDesc.attrs) + } catch let error { + throw error } - let input: Texture

-// let shape: [Int] - let inplace: Bool - var output: Texture

+ } + let input: Texture

+ // let shape: [Int] + let inplace: Bool + var output: Texture

} class ReshapeOp: Operator, ReshapeParam

>, Runable, Creator, InferShaperable{ - - func inferShape() { - // para.output.dim = para.input.dim - } - - typealias OpType = ReshapeOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } - } - func delogOutput() { - print("reshape delog") - let _: P? = para.input.metalTexture.logDesc(header: "reshape input: ", stridable: false) - let _: P? = para.output.metalTexture.logDesc(header: "reshape output: ", stridable: false) + + func inputs() -> [Variant] { + return [para.input] + } + + func inferShape() { + // para.output.dim = para.input.dim + } + + typealias OpType = ReshapeOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } + } + func delogOutput() { + print("reshape delog") +// let _: P? = para.input.metalTexture.logDesc(header: "reshape input: ", stridable: false) + let _: P? = para.output.metalTexture.logDesc(header: "reshape output: ", stridable: false) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/SoftmaxOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/SoftmaxOp.swift index d323b21cfa7729876a78702d0098c267132b4ab1..d212ba715b86858470161b9dff45e1454d3752f3 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/SoftmaxOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/SoftmaxOp.swift @@ -15,36 +15,40 @@ import Foundation class SoftmaxParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - input = try SoftmaxParam.inputX(inputs: opDesc.inputs, from: inScope) - output = try SoftmaxParam.outputOut(outputs: opDesc.outputs, from: inScope) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + input = try SoftmaxParam.inputX(inputs: opDesc.inputs, from: inScope) + output = try SoftmaxParam.outputOut(outputs: opDesc.outputs, from: inScope) + } catch let error { + throw error } - let input: Texture

- var output: Texture

+ } + let input: Texture

+ var output: Texture

} class SoftmaxOp: Operator, SoftmaxParam

>, Runable, Creator, InferShaperable{ - - func inferShape() { - // para.output.dim = para.input.dim - } - - typealias OpType = SoftmaxOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } - } - func delogOutput() { - print("softmax delog") - let _: P? = para.input.metalTexture.logDesc(header: "softmax input: ", stridable: false) - let _: P? = para.output.metalTexture.logDesc(header: "softmax output: ", stridable: false) + + func inputs() -> [Variant] { + return [para.input] + } + + func inferShape() { + // para.output.dim = para.input.dim + } + + typealias OpType = SoftmaxOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error } + } + func delogOutput() { + print("softmax delog") + let _: P? = para.input.metalTexture.logDesc(header: "softmax input: ", stridable: false) + let _: P? = para.output.metalTexture.logDesc(header: "softmax output: ", stridable: false) + } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift index 02edab912865f59181a3e9e7e3b93a4310806577..387353e236146c43a263951792f22c9827a02287 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift @@ -15,35 +15,54 @@ import Foundation class TransposeParam: OpParam { - typealias ParamPrecisionType = P - required init(opDesc: OpDesc, inScope: Scope) throws { - do { - input = try TransposeParam.inputX(inputs: opDesc.inputs, from: inScope) - output = try TransposeParam.outputOut(outputs: opDesc.outputs, from: inScope) - axis = try TransposeParam.getAttr(key: "axis", attrs: opDesc.attrs) - } catch let error { - throw error - } + typealias ParamPrecisionType = P + required init(opDesc: OpDesc, inScope: Scope) throws { + do { + input = try TransposeParam.inputX(inputs: opDesc.inputs, from: inScope) + output = try TransposeParam.outputOut(outputs: opDesc.outputs, from: inScope) + axis = try TransposeParam.getAttr(key: "axis", attrs: opDesc.attrs) + } catch let error { + throw error } - let input: Texture

- var output: Texture

- let axis: [Int32] + } + let input: Texture

+ var output: Texture

+ let axis: [Int32] } class TransposeOp: Operator, TransposeParam

>, Runable, Creator, InferShaperable{ - - func inferShape() { - para.output.dim = para.input.dim + + func inputs() -> [Variant] { + return [para.input] + } + + func inferShape() { + //para.output.dim = para.input.dim + } + + typealias OpType = TransposeOp

+ func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { + do { + try kernel.compute(commandBuffer: buffer, param: para) + } catch let error { + throw error + } + } + func delogOutput() { + let inputArray: [Float32] = para.input.metalTexture.floatArray { (ele: Float32) -> Float32 in + return ele } - typealias OpType = TransposeOp

- func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - do { - try kernel.compute(commandBuffer: buffer, param: para) - } catch let error { - throw error - } + print(inputArray.strideArray()) + + let outputArray: [Float32] = para.output.metalTexture.floatArray { (ele: Float32) -> Float32 in + return ele } + + print(outputArray.strideArray()) +// writeToLibrary(fileName: "transpose_ouput", array: outputArray) + } + } diff --git a/metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift b/metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift index be37b85e1f5ff71d7aaef38c7d6f8b8996967805..d17f56cead5a40aede28c6bb76d77b6384aae552 100644 --- a/metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift +++ b/metal/paddle-mobile/paddle-mobile/Program/TensorDesc.swift @@ -17,7 +17,7 @@ import Foundation struct TensorDesc { let dims: [Int] let dataType: VarTypeType - let dataLayout: DataLayout = DataLayout.NHWC() + let dataLayout: DataLayout = DataLayout.NCHW() var NCHWDim: [Int] { get { if dims.count != 4 { @@ -53,7 +53,7 @@ struct TensorDesc { } init(protoTensorDesc: PaddleMobile_Framework_Proto_VarType.TensorDesc) { - dims = protoTensorDesc.dims.map{ Int($0) > 0 ? Int($0) : 1 } + dims = protoTensorDesc.dims.map{ Int($0) > 0 ? Int($0) : abs(Int($0)) } dataType = VarTypeType.init(rawValue: protoTensorDesc.dataType.rawValue) ?? .ErrorType } diff --git a/metal/paddle-mobile/paddle-mobile/framework/Tensor.swift b/metal/paddle-mobile/paddle-mobile/framework/Tensor.swift index d2c8228ed9d9dd7e439aa1f688c60579f615b8a1..45e9eeb0009c2f3daeb11f6f879343d4790cc4e9 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Tensor.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Tensor.swift @@ -174,7 +174,7 @@ class Tensor: Tensorial { fatalError(" not support !") } //TODO: release - data.release() +// data.release() } var width: Int { diff --git a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift index fb95ee65f066c4f454781e2767106cc2f06d65fc..231ef0fb278d0e634c853d7a84a06b7fb9acedee 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift @@ -16,100 +16,103 @@ import Metal import Foundation class InputTexture { - let mtlTexture: MTLTexture - let expectDim: Dim - init(inMTLTexture: MTLTexture, inExpectDim: Dim) { - mtlTexture = inMTLTexture - expectDim = inExpectDim - } + let mtlTexture: MTLTexture + let expectDim: Dim + init(inMTLTexture: MTLTexture, inExpectDim: Dim) { + mtlTexture = inMTLTexture + expectDim = inExpectDim + } } extension InputTexture { - var description: String { - get{ - return mtlTexture.description - } + var description: String { + get{ + return mtlTexture.description } - - var debugDescription: String { - get { - return mtlTexture.debugDescription ?? " MetalTexture " - } + } + + var debugDescription: String { + get { + return mtlTexture.debugDescription ?? " MetalTexture " } + } } public class Texture: Tensorial { - var dim: Dim - var tensorDim: Dim - private(set) public var originDim: Dim - private var textureDesc: MTLTextureDescriptor! - public var metalTexture: MTLTexture! - var transpose: [Int] = [0, 1, 2, 3] + var dim: Dim + private(set) public var tensorDim: Dim + private(set) public var originDim: Dim + private var textureDesc: MTLTextureDescriptor! + public var metalTexture: MTLTexture! + var transpose: [Int] = [0, 1, 2, 3] + + func initTexture(device: MTLDevice, inTranspose: [Int] = [0, 1, 2, 3]) { + transpose = inTranspose + let newDim = transpose.map { originDim[$0] } - func initTexture(device: MTLDevice, transpose: [Int] = [0, 1, 2, 3]) { - 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() - - 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 = .rgba32Float - } - - tmpTextureDes.usage = [.shaderRead, .shaderWrite] - tmpTextureDes.storageMode = .shared - textureDesc = tmpTextureDes - metalTexture = device.makeTexture(descriptor: tmpTextureDes) ?! " texture nil " - } + let newLayout = transpose.map { layout.layoutWithDim[$0] } + + layout = DataLayout.init(newLayout) + dim = Dim.init(inDim: newDim) + + let tmpTextureDes = MTLTextureDescriptor.init() + + tmpTextureDes.width = newDim[2] + // layout.W ?? 1 + tmpTextureDes.height = newDim[1] + // layout.H ?? 1 + tmpTextureDes.depth = 1 + tmpTextureDes.arrayLength = ((newDim[0]) * (newDim[3]) + 3) / 4 + tmpTextureDes.textureType = .type2DArray - 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 ") - } - tensorDim = inDim - dim = fourDim - originDim = fourDim - layout = DataLayout.init([(.N, fourDim[0]), (.C, fourDim[1]), (.H, fourDim[2]), (.W, fourDim[3])]) + if MemoryLayout

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

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

.size == 4 { + tmpTextureDes.pixelFormat = .rgba32Float } - private(set) var layout: DataLayout + tmpTextureDes.usage = [.shaderRead, .shaderWrite] + tmpTextureDes.storageMode = .shared + textureDesc = tmpTextureDes + 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 ") + } + tensorDim = inDim + dim = fourDim + originDim = fourDim + layout = DataLayout.init([(.N, fourDim[0]), (.C, fourDim[1]), (.H, fourDim[2]), (.W, fourDim[3])]) + } + + private(set) var layout: DataLayout } extension Texture { - public var description: String { - return debugDescription - } - - public var debugDescription: String{ - var str = "" - str += "Dim: \(dim) \n value:[ " - str += "\(metalTexture)" - str += " ]" - return str - } - + public var description: String { + return debugDescription + } + + public var debugDescription: String{ + var str = "" + str += "Dim: \(dim) \n value:[ " + str += "\(metalTexture)" + str += " ]" + return str + } + }