diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/xcschememanagement.plist b/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/xcschememanagement.plist index e2c6b20680bd2f1209dfea198f9f62f2171716cf..125fd5ec745f635929df2d4e29f8147ab9a6b83a 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/xcschememanagement.plist +++ b/metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/xcschememanagement.plist @@ -7,7 +7,7 @@ paddle-mobile-demo.xcscheme orderHint - 3 + 4 diff --git a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift index 1396e95f4e61b85627e1d9349eb24e91be7e6d6a..7c06d16c145bc6c0affeb214ec54a400180bf101 100644 --- a/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift +++ b/metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift @@ -15,14 +15,41 @@ import UIKit import MetalKit import paddle_mobile +import MetalPerformanceShaders + + + +func Test() -> T? { + return nil +} class ViewController: UIViewController { let device: MTLDevice! = MTLCreateSystemDefaultDevice() var textureLoader: MTKTextureLoader! // let queue: MTLCommandQueue + func scaleTexture(queue: MTLCommandQueue, input: MTLTexture, complete: @escaping (MTLTexture) -> Void) { + let tmpTextureDes = MTLTextureDescriptor.init() + tmpTextureDes.width = 227 + tmpTextureDes.height = 227 + tmpTextureDes.depth = 1 + tmpTextureDes.usage = [.shaderRead, .shaderWrite] + tmpTextureDes.pixelFormat = .rgba16Float + tmpTextureDes.textureType = .type2D + tmpTextureDes.storageMode = .shared + tmpTextureDes.cpuCacheMode = .defaultCache + let dest = device.makeTexture(descriptor: tmpTextureDes) + + let scale = MPSImageLanczosScale.init(device: device) + let buffer = queue.makeCommandBuffer() + scale.encode(commandBuffer: buffer!, sourceTexture: input, destinationTexture: dest!) + buffer?.addCompletedHandler({ (buffer) in + complete(dest!) + }) + buffer?.commit() + } + override func viewDidLoad() { super.viewDidLoad() - let queue = device.makeCommandQueue() textureLoader = MTKTextureLoader.init(device: device) @@ -35,18 +62,24 @@ class ViewController: UIViewController { guard let inTexture = texture else { fatalError(" texture is nil !") } - - let loader = Loader.init() - do { - let modelPath = Bundle.main.path(forResource: "model", ofType: nil) ?! "model null" - let paraPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null" - let program = try loader.load(device: device, modelPath: modelPath, paraPath: paraPath) - let executor = try Executor.init(inDevice: device, inQueue: queue!, inProgram: program) - let output = try executor.predict(input: inTexture, expect: [1, 227, 227, 3]) - print(output) - } catch let error { - print(error) + + scaleTexture(queue: queue!, input: inTexture) { (inputTexture) in + let loader = Loader.init() + do { + let modelPath = Bundle.main.path(forResource: "model", ofType: nil) ?! "model null" + let paraPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null" + let program = try loader.load(device: self.device, modelPath: modelPath, paraPath: paraPath) + let executor = try Executor.init(inDevice: self.device, inQueue: queue!, inProgram: program) + let output = try executor.predict(input: inputTexture, expect: [1, 227, 227, 3]) + // print(output) + } catch let error { + print(error) + } } + + + + } } diff --git a/metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj b/metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj index 5234f24ddb107c5d809e873eafa8ad297e4cdb95..68023ae43a821b55893f8a3ac1c8b921d51855ed 100644 --- a/metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj +++ b/metal/paddle-mobile/paddle-mobile.xcodeproj/project.pbxproj @@ -36,6 +36,7 @@ FC0E2DC020EE461F009C1FAC /* ElementwiseAddKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */; }; FC1B16B320EC9A4F00678B91 /* Kernels.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC1B16B220EC9A4F00678B91 /* Kernels.metal */; }; FC1B186620ECF1C600678B91 /* ResizeKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC1B186520ECF1C600678B91 /* ResizeKernel.swift */; }; + FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */; }; FC60DB8920E9AAA500FF203F /* MetalExtension.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC60DB8820E9AAA500FF203F /* MetalExtension.swift */; }; FC82735920E3C04200BE430A /* OpCreator.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC82735820E3C04200BE430A /* OpCreator.swift */; }; FC9D037920E229E4000F735A /* OpParam.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037820E229E4000F735A /* OpParam.swift */; }; @@ -79,6 +80,7 @@ FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ElementwiseAddKernel.swift; sourceTree = ""; }; FC1B16B220EC9A4F00678B91 /* Kernels.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = Kernels.metal; sourceTree = ""; }; FC1B186520ECF1C600678B91 /* ResizeKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ResizeKernel.swift; sourceTree = ""; }; + FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Texture2DTo2DArrayKernel.swift; sourceTree = ""; }; FC60DB8820E9AAA500FF203F /* MetalExtension.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MetalExtension.swift; sourceTree = ""; }; FC82735820E3C04200BE430A /* OpCreator.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpCreator.swift; sourceTree = ""; }; FC9D037820E229E4000F735A /* OpParam.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpParam.swift; sourceTree = ""; }; @@ -212,6 +214,7 @@ FC0E2DB920EE3B8D009C1FAC /* ReluKernel.swift */, FC0E2DBD20EE460D009C1FAC /* BatchNormKernel.swift */, FC0E2DBF20EE461F009C1FAC /* ElementwiseAddKernel.swift */, + FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */, ); path = Kernels; sourceTree = ""; @@ -356,6 +359,7 @@ FC9D038220E2312E000F735A /* FetchOp.swift in Sources */, FC039BBD20E11CC20081E9F8 /* Program.swift in Sources */, FC039BA220E11CB70081E9F8 /* Loader.swift in Sources */, + FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */, FC039BC020E11CC20081E9F8 /* BlockDesc.swift in Sources */, FC039BAD20E11CBC0081E9F8 /* ReluOp.swift in Sources */, FC039BBE20E11CC20081E9F8 /* OpDesc.swift in Sources */, diff --git a/metal/paddle-mobile/paddle-mobile.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/xcschememanagement.plist b/metal/paddle-mobile/paddle-mobile.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/xcschememanagement.plist index 50f16e4d7cfc755905c68ced31769d3ac1dd1049..52725213607bd68fccbd4aa63faa9294cb622962 100644 --- a/metal/paddle-mobile/paddle-mobile.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/xcschememanagement.plist +++ b/metal/paddle-mobile/paddle-mobile.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/xcschememanagement.plist @@ -7,7 +7,7 @@ paddle-mobile.xcscheme orderHint - 4 + 3 diff --git a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift index 62dfdfa9f96e631f53cc093611217e7e9b5b92df..dbd015c3eb89288897f0464fd6101b827aacbeee 100644 --- a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift +++ b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift @@ -29,7 +29,6 @@ extension MTLDevice { fatalError("Counld't find paddle mobile library") } do { - print(path) paddleMobileMetalLibrary = try makeLibrary(filepath: path) } catch _ { fatalError("Counld't load paddle mobile library") @@ -61,22 +60,21 @@ extension MTLDevice { extension MTLComputeCommandEncoder { func dispatch(computePipline: MTLComputePipelineState, outTexture: MTLTexture) { - let slices = (outTexture.depth + 3)/4 + let slices = (outTexture.arrayLength * 4 + 3)/4 let width = computePipline.threadExecutionWidth let height = computePipline.maxTotalThreadsPerThreadgroup/width let threadsPerGroup = MTLSize.init(width: width, height: height, depth: 1) - print(" threads per group: \(threadsPerGroup) ") - - print(" out texture width: \(outTexture.width) , out texture height: \(outTexture.height)") +// print(" thread: threads per group: \(threadsPerGroup) ") +// print(" thread: out texture width: \(outTexture.width) , out texture height: \(outTexture.height)") let groupWidth = (outTexture.width + width - 1)/width let groupHeight = (outTexture.height + height - 1)/height let groupDepth = slices let groups = MTLSize.init(width: groupWidth, height: groupHeight, depth: groupDepth) - print("groups: \(groups) ") +// print("groups: \(groups) ") setComputePipelineState(computePipline) dispatchThreadgroups(groups, threadsPerThreadgroup: threadsPerGroup) @@ -84,6 +82,60 @@ extension MTLComputeCommandEncoder { } +public extension MTLTexture { + func logDesc(header: String = "", stridable: Bool = true) -> T? { + print(header) + print("texture: \(self)") + 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 { + for j in stride(from: 0, to: width * height * depth * 4 , by: width * height * depth * 4 / 100){ + str += " \(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 += " \(p[j])" + } + } else { + for j in 0.. { } buffer.addCompletedHandler { (commandbuffer) in + +// for op in self.ops { +// op.delogOutput() +// } + let afterDate = Date.init() - print(afterDate.timeIntervalSince(beforeDate)) - print(" encoder end ! ") + print(" encoder end ! time: \(afterDate.timeIntervalSince(beforeDate))") + } buffer.commit() diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift b/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift index 222c5319c6d7dfd88d8a83140bef98e67dcb2fd5..5253f2c4a554fb45d9075cee344160788c9c935b 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Base/Operator.swift @@ -18,6 +18,7 @@ import Foundation protocol Runable { func run(device: MTLDevice, buffer: MTLCommandBuffer) throws func runImpl(device: MTLDevice,buffer: MTLCommandBuffer) throws + func delogOutput() } extension Runable where Self: OperatorProtocol{ @@ -27,8 +28,11 @@ extension Runable where Self: OperatorProtocol{ } catch let error { throw error } - - print(type + ": " + para.outputDesc()) +// print(type + ": " + para.outputDesc()) + } + + func delogOutput() { + print(type + ": has no implementation" ) } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift index 150004f2ce8c718092d255226b8e1c497fc4bba0..4f0a12c5fd2e1e6e452eb56c10b598565f30c499 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/BatchNormOp.swift @@ -48,7 +48,6 @@ class BatchNormOp: Operator, BatchNormKernel } typealias OpType = BatchNormOp

func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - print("this is BatchNormOp") } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift index b4623bf567cbcb17c99b28e246d3bf0d242bc23b..aea34976fcc2e8d1518966055112ed1509598c82 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConvOp.swift @@ -64,6 +64,17 @@ class ConvOp: Operator, ConvKernel

>, Runable, typealias OpType = ConvOp

func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - print("this is conv") + 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/ElementwiseAddOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ElementwiseAddOp.swift index 12c0b832b76b0975d4b5a138c283866166e06130..18aed606f9be9d2e03d1927ffdae8c4144e5bc51 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ElementwiseAddOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ElementwiseAddOp.swift @@ -40,7 +40,6 @@ class ElementwiseAddOp: Operator, Eleme typealias OpType = ElementwiseAddOp

func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - print("this is ElementwiseAddOp") } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift index 87b71a7a95d316cd226e4e648506fbe011d5e4c8..eea42425423f391771338213a7d34ca3a42decbe 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/FeedOp.swift @@ -33,26 +33,37 @@ struct FeedParam: OpParam{ typealias ParamPrecisionType = P } -class FeedOp: Operator, ResizeKernel

>, Runable, Creator, InferShaperable { +class FeedOp: Operator, Texture2DTo2DArrayKernel

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

func inferShape() { -// print("feed input: \(para.input.expectDim)") + // print("feed input: \(para.input.expectDim)") print("feed output: \(para.output.dim)") // para.output.dim = -// para.output.dim = para.input.expectDim + // para.output.dim = para.input.expectDim } func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - let resizeKernel = ResizeKernel

.init(device: device) - let resizeParam = ResizeParam.init(input: para.input.mtlTexture, output: para.output.metalTexture, expectDim: para.input.expectDim) + let locPara = Texture2DTo2DArrayParam.init(input: para.input.mtlTexture, output: para.output.metalTexture, expectDim: para.input.expectDim) do { - print("feed op to compute ") - try resizeKernel.compute(commandBuffer: buffer, param: resizeParam) - print("feed op end compute ") + try kernel.compute(commandBuffer: buffer, param: locPara) } 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 _: Float16? = para.input.mtlTexture.logDesc(header: "feed input: ") + let _: Float16? = para.output.metalTexture.logDesc(header: "feed output: ") } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift index b79538fe7770dd818eb7399547d2621249938c89..a671c0fa45f9f7d4a31b31e8caa816900292da71 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/FetchOp.swift @@ -32,13 +32,11 @@ struct FetchParam: OpParam{ class FetchOp: Operator, ResizeKernel

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

func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws { - print("fetch op") } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvKernel.swift index 0cdab6f6b967214a0d9e8a11ae8768b252ead38e..e153233d0becb40354bfa63aee1049b0c5b8182e 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvKernel.swift @@ -11,10 +11,18 @@ import Foundation class ConvKernel: Kernel, Computable { func compute(commandBuffer: MTLCommandBuffer, param: ConvParam

) 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.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) + encoder.endEncoding() } required init(device: MTLDevice) { super.init(device: device, inFunctionName: "conv") } + + } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal index 0d872196c9d19497fa37a35dc9b294843624ca56..56d9ca78ce93ac3f5159b7d415db72c42d8f751c 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal @@ -36,7 +36,6 @@ kernel void resize(texture2d inTexture [[texture(0)]], outTexture.write(half4(input.x, input.y, input.z, input.w), gid.xy, gid.z); } - kernel void relu(texture2d_array inTexture [[texture(0)]], texture2d_array outTexture [[texture(1)]], uint3 gid [[thread_position_in_grid]]) { @@ -49,7 +48,6 @@ kernel void relu(texture2d_array inTexture [[texture(0)]], outTexture.write(half4(relu), gid.xy, gid.z); } - kernel void elementwise_add(texture2d_array inTexture [[texture(0)]], texture2d_array outTexture [[texture(1)]], const device half4 *biasTerms [[buffer(0)]], @@ -62,10 +60,8 @@ kernel void elementwise_add(texture2d_array inTexture [[text outTexture.write(input, gid.xy, gid.z); } - kernel void conv(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() || @@ -75,17 +71,27 @@ kernel void conv(texture2d_array inTexture [[texture(0)]], outTexture.write(input, gid.xy, gid.z); } + kernel void batchnorm(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; - 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); } +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 half4 input = inTexture.read(gid.xy); + outTexture.write(input, gid.xy, 0); +} diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ReluKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ReluKernel.swift index e14cf9f942d01a720b44e4736d176768944947cd..614b47dc8c70fb7e63371873a273c67ac81be749 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ReluKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ReluKernel.swift @@ -19,7 +19,6 @@ class ReluKernel: Kernel, Computable{ guard let encoder = commandBuffer.makeComputeCommandEncoder() else { throw PaddleMobileError.predictError(message: " encode is nil") } - print(" the usage of input of relu \(param.input.metalTexture.usage)") encoder.setTexture(param.input.metalTexture, index: 0) encoder.setTexture(param.output.metalTexture, index: 1) encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture) diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ResizeKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ResizeKernel.swift index 9de09ccb98d71eb41ae576ad19677b7f46eea70b..f938c1379585f529b8b96cba20a628826dea7435 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ResizeKernel.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ResizeKernel.swift @@ -13,6 +13,7 @@ limitations under the License. */ import Foundation +import MetalPerformanceShaders struct ResizeParam { @@ -29,23 +30,29 @@ struct OutputDim { } class ResizeKernel: Kernel, Computable{ + var lanczos: MPSImageLanczosScale + required init(device: MTLDevice) { + lanczos = MPSImageLanczosScale.init(device: device) + super.init(device: device, inFunctionName: "resize") + } func compute(commandBuffer: MTLCommandBuffer, param: ResizeParam) throws { - guard let encoder = commandBuffer.makeComputeCommandEncoder() else { - throw PaddleMobileError.predictError(message: " encode is nil") - } +// guard let encoder = commandBuffer.makeComputeCommandEncoder() else { +// throw PaddleMobileError.predictError(message: " encode is nil") +// } + lanczos.encode(commandBuffer: commandBuffer, sourceTexture: param.input, destinationTexture: param.output) - encoder.setTexture(param.input, index: 0) - encoder.setTexture(param.output, index: 1) - let strideX = param.input.width/param.expectDim[2] - let strideY = param.input.height/param.expectDim[1] - var outputDim = OutputDim.init(width: UInt16(param.expectDim[1]), height: UInt16(param.expectDim[2]), strideX: UInt16(strideX), strideY: UInt16(strideY)) - encoder.setBytes(&outputDim, length: MemoryLayout.size, index: 0) - encoder.dispatch(computePipline: pipline, outTexture: param.output) - encoder.endEncoding() +// encoder.setTexture(param.input, index: 0) +// encoder.setTexture(param.output, index: 1) +// let strideX = param.input.width/param.expectDim[2] +// let strideY = param.input.height/param.expectDim[1] +// var outputDim = OutputDim.init(width: UInt16(param.expectDim[1]), height: UInt16(param.expectDim[2]), strideX: UInt16(strideX), strideY: UInt16(strideY)) +// encoder.setBytes(&outputDim, length: MemoryLayout.size, index: 0) +// encoder.dispatch(computePipline: pipline, outTexture: param.output) +// encoder.endEncoding() } - required init(device: MTLDevice) { - super.init(device: device, inFunctionName: "resize") - } + + + } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Texture2DTo2DArrayKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Texture2DTo2DArrayKernel.swift new file mode 100644 index 0000000000000000000000000000000000000000..ea471ae13d296a18a0f08c9b0e820a0666084a99 --- /dev/null +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Texture2DTo2DArrayKernel.swift @@ -0,0 +1,33 @@ +// +// Texture2DTo2DArrayKernel.swift +// paddle-mobile +// +// Created by liuRuiLong on 2018/7/6. +// Copyright © 2018年 orange. All rights reserved. +// + +import Foundation + +struct Texture2DTo2DArrayParam { + let input: MTLTexture + let output: MTLTexture + let expectDim: Dim +} + + +class Texture2DTo2DArrayKernel: Kernel, Computable{ + func compute(commandBuffer: MTLCommandBuffer, param: Texture2DTo2DArrayParam) throws { + guard let encoder = commandBuffer.makeComputeCommandEncoder() else { + throw PaddleMobileError.predictError(message: " encode is nil") + } + encoder.setTexture(param.input, index: 0) + encoder.setTexture(param.output, index: 1) + encoder.dispatch(computePipline: pipline, outTexture: param.input) + encoder.endEncoding() + } + + required init(device: MTLDevice) { + super.init(device: device, inFunctionName: "texture2d_to_2d_array") + } +} + diff --git a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift index 141dc70df2e966516eaa340cbecbaa6ccb722697..9ea5c38ba1de351425c9375043d0432fba731be2 100644 --- a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift +++ b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift @@ -62,14 +62,16 @@ public class Texture: Tensorial { fatalError(" didn't support yet") } if MemoryLayout

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

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

.size == 4 { - tmpTextureDes.pixelFormat = .r32Float +// tmpTextureDes.pixelFormat = .r32Float + tmpTextureDes.pixelFormat = .rgba32Float + } - tmpTextureDes.usage = .unknown + tmpTextureDes.usage = [.shaderRead, .shaderWrite] tmpTextureDes.storageMode = .shared textureDesc = tmpTextureDes metalTexture = device.makeTexture(descriptor: tmpTextureDes) ?! " texture nil " @@ -123,6 +125,7 @@ extension Texture { public var debugDescription: String{ var str = "" str += "Dim: \(dim) \n value:[ " +// str += "\(metalTexture)" str += " ]" return str }