diff --git a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift index 6e01c4fb27163d52d742f219ee80fe76ed6fc8ca..14e492e761feece6fdfd5554a5ca584c81a176a3 100644 --- a/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift +++ b/metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift @@ -381,7 +381,7 @@ public extension MTLTexture { for w in 0.. dim.c { - for i in 0..<((sliceIndex * 4 + 4) - dim.c) { + for i in 0..<(4 - ((sliceIndex * 4 + 4) - dim.c)) { let value = textureArray[sliceIndex * numOfASlice + h * dim.w * 4 + w * 4 + i] output.append(value) } diff --git a/metal/paddle-mobile/paddle-mobile/Executor.swift b/metal/paddle-mobile/paddle-mobile/Executor.swift index d9e2ae9fb1d4a66130dfb37e64eb70f02a833250..dcb4a72bd886177f1cc954f15128d1be982ddf24 100644 --- a/metal/paddle-mobile/paddle-mobile/Executor.swift +++ b/metal/paddle-mobile/paddle-mobile/Executor.swift @@ -14,7 +14,7 @@ import Foundation -let testTo = 48 +let testTo = 93 public class ResultHolder { public let dim: [Int] diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift index 7c1532adf9fbe75ba842ab12fb74d73f90f0899c..72711bcff02ea70f9041edaeb37a2e8d098d428e 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift @@ -59,15 +59,21 @@ class ConcatOp: Operator, ConcatParam

>, Run } func delogOutput() { - let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in - return o - } -// print(outputArray.strideArray()) + let originDim = para.output.originDim + let outputArray = para.output.metalTexture.realNHWC(dim: (n: originDim[0], h: originDim[1], w: originDim[2], c: originDim[3])) + print(outputArray.strideArray()) - writeToLibrary(fileName: "concat_out", array: outputArray) - let device: MTLDevice = MTLCreateSystemDefaultDevice()! +// let outputArray = para.output.metalTexture.floatArray { (o: Float32) -> Float32 in +// return o +// } +// +//// print(outputArray.strideArray()) +// +// writeToLibrary(fileName: "concat_out", array: outputArray) +// +// let device: MTLDevice = MTLCreateSystemDefaultDevice()! // let tensorArray: [P] = device.texture2tensor(texture: para.output.metalTexture, dim: [1917, 4]) diff --git a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal index b4468af42838d028270e5f62c320ddfeedddee98..00ba9c164668c2283c0f47320d9fb25c574c3b36 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal @@ -24,6 +24,52 @@ struct ReshapeParam { 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(texture2d_array inTexture [[texture(0)]], texture2d_array outTexture [[texture(1)]], constant ReshapeParam &rp [[buffer(0)]], @@ -32,7 +78,7 @@ kernel void reshape(texture2d_array inTexture [[texture(0)] 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]; + 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]]; @@ -40,26 +86,15 @@ kernel void reshape(texture2d_array inTexture [[texture(0)] 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]; + index2abcd(lrp.idim, index, tabcd); + trans(lrp.itrans, tabcd, iabcd); + abcd2xyzn(iC, iabcd, ixyzn); + r[n] = inTexture.read(uint2(ixyzn[0], ixyzn[1]), ixyzn[2])[ixyzn[3]]; } else { r[n] = 0; } @@ -67,42 +102,6 @@ kernel void reshape(texture2d_array inTexture [[texture(0)] outTexture.write(r, gid.xy, gid.z); } - -/* - - 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]; - 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; - } - } - outTexture.write(r, gid.xy, gid.z); - } - */ - // //kernel void reshape_half(texture2d_array inTexture [[texture(0)]], // texture2d_array outTexture [[texture(1)]], diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift index d637a176bcda9878b542c39ab6704bc0240a7149..3bcf9c15a0fe2d075e79ee98f2f32726efbdbad7 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift @@ -53,10 +53,11 @@ class ReshapeOp: Operator, ReshapeParam

>, 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: true) - let deivice = MTLCreateSystemDefaultDevice()! - let array: [Float32] = deivice.texture2tensor(texture: para.output.metalTexture, dim: [1, 1, 600, 7]) - print(array.strideArray()) + let originDim = para.output.originDim + + let outputArray = para.output.metalTexture.realNHWC(dim: (n: originDim[0], h: originDim[1], w: originDim[2], c: originDim[3])) + print(outputArray.strideArray()) + } }