diff --git a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift
index b750018260f64ae89f5b3aab5cc987eee9a11415..71b50f0553cbef8d1f0815f95c4ce9e4af150021 100644
--- a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift
+++ b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift
@@ -71,6 +71,105 @@ extension MTLDevice {
return buffer!
}
+ 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], 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()
@@ -85,19 +184,25 @@ extension MTLDevice {
textureDesc.arrayLength = arrayLength
let texture = makeTexture(descriptor: textureDesc)!
- if arrayLength == 1 && value.count >= 4{
- let pointer: UnsafeMutablePointer = UnsafeMutablePointer
.allocate(capacity: value.count * MemoryLayout
.size)
+ 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))
- texture.replace(region: region, mipmapLevel: 0, withBytes: pointer, bytesPerRow: bytesPerRow)
+ for i in 0..(header: String = "", stridable: Bool = true) -> T? {
print(header)
print("texture: \(self)")
- let res: [(index: Int, value: T)] = stridableFloatArray(stridable: stridable)
- print(res)
+// 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 > 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...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.. 0 {
+ log += ", "
+ }
+ log += tensor[c].description
+ }
+ log += "]"
+ if (indentLevel > 0) && (ix[indentLevel - 1] < dim[indentLevel - 1] - 1) {
+ log += ","
+ }
+ print(log)
+ } else {
+ print(indent + "[")
+ for i in 0.. 0) && (ix[indentLevel - 1] < dim[indentLevel - 1] - 1) {
+ print(indent + "],")
+ } else {
+ print(indent + "]")
+ }
+ }
+ }
+
+ private func tensorPrint(tensor: [Float32], dim: [Int]) {
+ var detectPos = -1
+ var odim = 1
+ var ndim = dim
+ for i in 0..= -1)
+ if (detectPos == -1) {
+ assert(tensor.count == odim)
+ } else {
+ assert(tensor.count % odim == 0)
+ ndim[detectPos] = tensor.count / odim
+ }
+ indentPrintTensor(tensor: tensor, dim: ndim, ix: dim.map { $0 * 0 }, indentLevel: 0)
+ }
+
+ public func testTranspose() {
+
+ let buffer = queue.makeCommandBuffer() ?! "buffer is nil"
+ var input: [Float32] = []
+ for i in 0..<72 {
+ input.append(Float32(i))
+ }
+// let inputTexture = device.makeFloatTexture(value: input, textureWidth: 3, textureHeight: 2, arrayLength: 3)
+ let inputTexture = device.tensor2texture(value: input, dim: [4, 3, 2, 3]);
+ // group 1
+ let outputTexture = device.tensor2texture(value: [Float32](), dim: [3, 3, 2, 4])
+ let param = TransposeTestParam.init(inputTexture: inputTexture, outputTexture: outputTexture, iC: 3, oC: 4, axis: [3, 1, 2, 0])
+// let param = TransposeTestParam.init(inputTexture: inputTexture, outputTexture: outputTexture, iC: 4, oC: 2, axis: [3, 0, 2, 1])
+// // group 2
+// let outputTexture = device.makeFloatTexture(value: [Float32](), textureWidth: 3, textureHeight: 3, arrayLength: 6)
+// let param = TransposeTestParam.init(inputTexture: inputTexture, outputTexture: outputTexture, iC: 4, oC: 4, axis: [3, 0, 2, 1])
+//
+ let transposeKernel = TransposeKernel.init(device: device, testParam: param)
+
+ transposeKernel.test(commandBuffer: buffer, param: param)
+
+ buffer.addCompletedHandler { (buffer) in
+ let _: Float32? = inputTexture.logDesc(header: "input texture", stridable: false)
+ let _: Float32? = outputTexture.logDesc(header: "output texture", stridable: false)
+ self.tensorPrint(tensor: input, dim: [4, 3, 2, 3])
+ let tx: [Float32] = self.device.texture2tensor(texture: outputTexture, dim: [3, 3, 2, 4])
+ self.tensorPrint(tensor: tx, dim: [3, 3, 2, 4])
+ }
+
+ buffer.commit()
+ }
+
public func testConvAddBnRelu() {
let buffer = queue.makeCommandBuffer() ?! " buffer is nil "
@@ -132,16 +229,6 @@ public class PaddleMobileUnitTest {
}
buffer.commit()
-
-
-// let inputTexture = device.makeFloatTexture(value: <#T##[P]#>, textureWidth: <#T##Int#>, textureHeight: <#T##Int#>, arrayLength: <#T##Int#>)
-
-
-// let param = ConvAddBatchNormReluTestParam.init(inInputTexture: <#T##MTLTexture#>, inOutputTexture: <#T##MTLTexture#>, inMetalParam: <#T##MetalConvParam#>, inFilterBuffer: <#T##MTLBuffer#>, inBiaseBuffer: <#T##MTLBuffer#>, inNewScaleBuffer: <#T##MTLBuffer#>, inNewBiaseBuffer: <#T##MTLBuffer#>, inFilterSize: <#T##(width: Int, height: Int, channel: Int)#>)
-
-// ConvAddBatchNormReluKernel.init(device: <#T##MTLDevice#>, testParam: <#T##ConvAddBatchNormReluTestParam#>)
-
-
}
}
diff --git a/metal/paddle-mobile/paddle-mobile/Executor.swift b/metal/paddle-mobile/paddle-mobile/Executor.swift
index 0dcb3151e21cc0f3968a07da39366d4ba5fd5813..fc19f32ebc588fcb2f9dab7b0f92ad0d28f2efa4 100644
--- a/metal/paddle-mobile/paddle-mobile/Executor.swift
+++ b/metal/paddle-mobile/paddle-mobile/Executor.swift
@@ -68,16 +68,6 @@ public class Executor {
throw error
}
}
-
-// for op in block.ops {
-// do {
-// let op = try OpCreator.shared.creat(device: inDevice, opDesc: op, scope: inProgram.scope)
-// op.inferShape()
-// ops.append(op)
-// } catch let error {
-// throw error
-// }
-// }
}
}
diff --git a/metal/paddle-mobile/paddle-mobile/Loader.swift b/metal/paddle-mobile/paddle-mobile/Loader.swift
index 2f9060a19f8ea67704711dd250e629e9c34f3eef..5c2733a5d9088d98a6c25f3efd040be5bc482d0d 100644
--- a/metal/paddle-mobile/paddle-mobile/Loader.swift
+++ b/metal/paddle-mobile/paddle-mobile/Loader.swift
@@ -141,10 +141,6 @@ public class Loader {
throw PaddleMobileError.loaderError(message: "get tensor desc failed")
}
-// guard (try? tensorDesc.dataType.dataTypeSize()) == MemoryLayout.size else {
-// throw PaddleMobileError.memoryError(message: "PrecisionType not support")
-// }
-
if (varDesc.persistable
&& varDesc.type != .FeedMiniBatch
&& varDesc.type != .FetchList) {
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Base/OpCreator.swift b/metal/paddle-mobile/paddle-mobile/Operators/Base/OpCreator.swift
index 7e8c9c1a5adc5b7934069526989ca58bee6cc27c..248a0162e4425ed13efcee46d4177a2600534524 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/Base/OpCreator.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/Base/OpCreator.swift
@@ -56,7 +56,9 @@ class OpCreator {
gBoxcoderType : BoxcoderOp.creat,
gConvBnReluType : ConvBNReluOp
.creat,
gDwConvBnReluType : DwConvBNReluOp
.creat,
- gMulticlassNMSType : MulticlassNMSOp
.creat]
+ gMulticlassNMSType : MulticlassNMSOp
.creat,
+ gTransposeType : TransposeOp
.creat,
+ gPriorBoxType : PriorBoxOp
.creat]
private init(){}
}
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift
index 8eeae8867ec6f40b0dadaf5ba96aa847516f368f..4cea0c7c58b8ff89196d49962eca2624d1fed6a1 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/BoxcoderOp.swift
@@ -18,19 +18,28 @@ class BoxcoderParam: OpParam {
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
- fatalError()
+ 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
}
}
- let input: Texture
+ let priorBox: Texture
+ let priorBoxVar: Texture
+ let targetBox: Texture
var output: Texture
+ let codeType: String
+ let boxNormalized: Bool
}
class BoxcoderOp: Operator, BoxcoderParam>, Runable, Creator, InferShaperable{
func inferShape() {
- para.output.dim = para.input.dim
+// para.output.dim = para.input.dim
}
typealias OpType = BoxcoderOp
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift
index a85fec5880738791abc3551f5b712a736ba1f46d..bf2a60c613b00f09733dcff644d8417fb073b20b 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift
@@ -18,19 +18,31 @@ class ConcatParam: OpParam {
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
- fatalError()
+ 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
}
}
- let input: Texture
+ var input: [Texture
] = []
var output: Texture
+ let axis: Int
}
class ConcatOp: Operator, ConcatParam>, Runable, Creator, InferShaperable{
func inferShape() {
- para.output.dim = para.input.dim
+ 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
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/BoxcoderKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/BoxcoderKernel.swift
index 6cc15628cb8bddf28730d88408358d37a94bed93..461652e8c3e53114b36b89a4be785443d3756dcc 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/BoxcoderKernel.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/BoxcoderKernel.swift
@@ -19,7 +19,7 @@ class BoxcoderKernel: Kernel, Computable{
guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
throw PaddleMobileError.predictError(message: " encode is nil")
}
- encoder.setTexture(param.input.metalTexture, index: 0)
+// encoder.setTexture(param.input.metalTexture, index: 0)
encoder.setTexture(param.output.metalTexture, index: 1)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Concat.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Concat.swift
new file mode 100644
index 0000000000000000000000000000000000000000..25f0a21bfff420566d06a59dca626805dd0ce6e0
--- /dev/null
+++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Concat.swift
@@ -0,0 +1,31 @@
+/* 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 ConcatKernel: Kernel, Computable{
+ func compute(commandBuffer: MTLCommandBuffer, param: ConcatParam) throws {
+ guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
+ throw PaddleMobileError.predictError(message: " encoder 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, param: ConcatParam
) {
+ super.init(device: device, inFunctionName: "concat")
+ }
+}
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConcatKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConcatKernel.swift
index c1cc11d16f28a5d1b91fc2ce9670b04492f25ec8..f5a16141fa77a1a03eedbdb5748bd321f667d226 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConcatKernel.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConcatKernel.swift
@@ -19,13 +19,13 @@ class ConcatKernel: Kernel, Computable{
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.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, param: ConcatParam) {
- super.init(device: device, inFunctionName: "priorbox")
+ super.init(device: device, inFunctionName: "concat")
}
}
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift
index 867787b10cf5527e4f52d2f904b8cbc1b7f626f8..5b8726f5596fdf19d4ae1186110d182c78ac0900 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/ConvBNReluKernel.swift
@@ -57,7 +57,7 @@ class ConvBNReluKernel: Kernel, Computable, Testable {
} else {
super.init(device: device, inFunctionName: "conv_add_batch_norm_relu_3x3")
}
-
+ param.output.initTexture(device: device, transpose: [0, 2, 3, 1])
param.filter.initBuffer(device: device, precision: Tensor.BufferPrecision.Float32)
param.variance.initBuffer(device: device)
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal
index 668ed443be7b4bdf13a9bb118f8ba3bd0f0c2dd7..851dd35faf670282a06f5cf06282efa0bed9cde1 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal
+++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/Kernels.metal
@@ -202,7 +202,7 @@ kernel void reshape_half(texture2d_array inTexture [[texture
gid.y >= outTexture.get_height() ||
gid.z >= outTexture.get_array_size()) return;
- half4 r = inTexture.read(uint2(0, 0), gid.z);
+ half4 r = inTexture.read(uint2(0, 0), gid.x);
outTexture.write(r, gid.xy, gid.z);
}
@@ -321,8 +321,53 @@ kernel void prior_box(texture2d_array inTexture [[texture(0
}
}
+void xyzn2abcd(uint C, uint xyzn[4], uint abcd[4]) {
+ abcd[1] = xyzn[0];
+ abcd[2] = xyzn[1];
+ uint t = xyzn[2] * 4 + xyzn[3];
+ abcd[0] = t / C;
+ abcd[3] = t % C;
+ return;
+}
+void abcd2xyzn(uint C, uint abcd[4], uint xyzn[4]) {
+ xyzn[0] = abcd[1];
+ xyzn[1] = abcd[2];
+ uint t = abcd[0] * C + abcd[3];
+ xyzn[2] = t / 4;
+ xyzn[3] = t % 4;
+ return;
+}
+struct TransposeParam {
+ 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 (uint i = 0; i < 4; i++) {
+ uint ixyzn[] = {gid.x, gid.y, gid.z, i};
+ uint 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[i] = rt[oxyzn[3]];
+ }
+ outTexture.write(r, gid.xy, gid.z);
+ }
+}
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/MulticlassNMSKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/MulticlassNMSKernel.swift
index a85931af5e6ef042af62f754181de2beccb740bb..7db6bcca8b013bc8396b879e52ec68a3f2abf20b 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/MulticlassNMSKernel.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/MulticlassNMSKernel.swift
@@ -19,8 +19,8 @@ class MulticlassNMSKernel: Kernel, Computable{
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.input.metalTexture, index: 0)
+// encoder.setTexture(param.output.metalTexture, index: 1)
encoder.dispatch(computePipline: pipline, outTexture: param.output.metalTexture)
encoder.endEncoding()
}
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift
index c3868dd0708a9859f3133c7b71461a64f194c079..d6bdc06f59417af088a9913a12fc792ae376d942 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/TransposeKernel.swift
@@ -14,18 +14,73 @@
import Foundation
-class TransposeKernel: Kernel, Computable{
+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]))
+ }
+}
+
+struct TransposeTestParam: TestParam {
+ 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
+ }
+ let realAxis = param.axis.map {invT[$0]}
+ var tmp = TransposeMetalParam.init(realAxis)
+ tmp.iC = Int32(param.input.dim[param.input.transpose[3]])
+ tmp.oC = Int32(param.output.dim[3])
+
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()
}
required init(device: MTLDevice, param: TransposeParam) {
- super.init(device: device, inFunctionName: "priorbox")
+ param.output.initTexture(device: device, transpose: [0, 1, 2, 3])
+ super.init(device: device, inFunctionName: "transpose")
}
-}
+ required init(device: MTLDevice, testParam: TransposeTestParam) {
+ super.init(device: device, inFunctionName: "transpose")
+ }
+
+ 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()
+ }}
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift
index dae584d52c6d2335908d28741c4171c8699eff80..7931d3b5611546e3e9f06fa556c061fa85fe6804 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/MulticlassNMSOp.swift
@@ -18,19 +18,22 @@ class MulticlassNMSParam: OpParam {
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
- fatalError()
+ scores = try MulticlassNMSParam.getFirstTensor(key: "Scores", map: opDesc.inputs, from: inScope)
+ bboxes = try MulticlassNMSParam.getFirstTensor(key: "BBoxes", map: opDesc.inputs, from: inScope)
+ output = try MulticlassNMSParam.outputOut(outputs: opDesc.outputs, from: inScope)
} catch let error {
throw error
}
}
- let input: Texture
+ let scores: Texture
+ let bboxes: Texture
var output: Texture
}
class MulticlassNMSOp: Operator, MulticlassNMSParam>, Runable, Creator, InferShaperable{
func inferShape() {
- para.output.dim = para.input.dim
+// para.output.dim = para.input.dim
}
typealias OpType = MulticlassNMSOp
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift
index 772ed0e2f8a558eb8d54c1eacc0dd695c82511f2..da7a60ba5b64ba8b07f3eecfde3ce99f9a31f861 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/PriorBoxOp.swift
@@ -18,13 +18,16 @@ class PriorBoxParam: OpParam {
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
- fatalError()
+ input = try PriorBoxParam.input(inputs: opDesc.inputs, from: inScope)
+ output = try PriorBoxParam.getFirstTensor(key: "Boxes", map: opDesc.outputs, from: inScope)
+ variances = try PriorBoxParam.getFirstTensor(key: "Variances", map: opDesc.outputs, from: inScope)
} catch let error {
throw error
}
}
let input: Texture
var output: Texture
+ let variances: Texture
}
class PriorBoxOp: Operator, PriorBoxParam>, Runable, Creator, InferShaperable{
@@ -36,7 +39,7 @@ class PriorBoxOp: Operator, PriorBoxParam
typealias OpType = PriorBoxOp
func runImpl(device: MTLDevice, buffer: MTLCommandBuffer) throws {
do {
- try kernel.compute(commandBuffer: buffer, param: para)
+ // try kernel.compute(commandBuffer: buffer, param: para)
} catch let error {
throw error
}
diff --git a/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift
index f9a249484acb8533b66ed56635bc8e728c9ba3dd..02edab912865f59181a3e9e7e3b93a4310806577 100644
--- a/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift
+++ b/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift
@@ -18,13 +18,16 @@ class TransposeParam: OpParam {
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
- fatalError()
+ 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]
}
class TransposeOp: Operator, TransposeParam>, Runable, Creator, InferShaperable{
diff --git a/metal/paddle-mobile/paddle-mobile/framework/Dim.swift b/metal/paddle-mobile/paddle-mobile/framework/Dim.swift
index 672484cd9d055bbe65a61d41017199dd79d6cdb2..d584744aae4b3d135cfb83ab51c7b158d77db45a 100644
--- a/metal/paddle-mobile/paddle-mobile/framework/Dim.swift
+++ b/metal/paddle-mobile/paddle-mobile/framework/Dim.swift
@@ -39,7 +39,6 @@ public struct Dim {
return dims[index];
}
-
private(set) var dims: [Int]
private init(){
fatalError()
diff --git a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift
index c26994116dc3ba5b22c41d4c377edd1de6876e9d..3e8cdf9328e0046a6d2b50131ff20dcafe056b4b 100644
--- a/metal/paddle-mobile/paddle-mobile/framework/Texture.swift
+++ b/metal/paddle-mobile/paddle-mobile/framework/Texture.swift
@@ -95,43 +95,6 @@ public class Texture: Tensorial {
layout = DataLayout.init([(.N, fourDim[0]), (.C, fourDim[1]), (.H, fourDim[2]), (.W, fourDim[3])])
}
-// required public init(inDim: Dim, inLayout: DataLayout = .NHWC, inTexture: MTLTexture) {
-// dim = inDim
-// layout = inLayout
-// metalTexture = inTexture
-// let tmpTextureDes = MTLTextureDescriptor.init()
-//
-// if inDim.cout() == 1 {
-// tmpTextureDes.width = inDim[0]
-// tmpTextureDes.textureType = .type1D
-// } else if inDim.cout() == 2 {
-// tmpTextureDes.height = inDim[0]
-// tmpTextureDes.width = inDim[1]
-// tmpTextureDes.textureType = .type2D
-// } else if inDim.cout() == 3 {
-// fatalError(" not support texture dim 3")
-// } else if inDim.cout() == 4 {
-// tmpTextureDes.height = inDim[1]
-// tmpTextureDes.width = inDim[2]
-// tmpTextureDes.depth = inDim[3] * inDim[1]
-// tmpTextureDes.textureType = .type2DArray
-// }
-//
-// tmpTextureDes.pixelFormat = .r32Float
-// tmpTextureDes.storageMode = .shared
-// textureDesc = tmpTextureDes
-// let device = MTLCreateSystemDefaultDevice()
-// metalTexture = device!.makeTexture(descriptor: tmpTextureDes)!
-// }
-
-// init() {
-// dim = Dim.init(inDim: [])
-// layout = .NCHW
-// let device = MTLCreateSystemDefaultDevice()
-// textureDesc = MTLTextureDescriptor.init()
-// metalTexture = device!.makeTexture(descriptor: textureDesc)!
-// }
-
private(set) var layout: DataLayout
}