diff --git a/metal/paddle-mobile/paddle-mobile/Executor.swift b/metal/paddle-mobile/paddle-mobile/Executor.swift index 9cee599209b8858f76cf13ac1cdd0ae78d2b1bc8..869d83729daa0ffac45a4083aa02af8a74425b13 100644 --- a/metal/paddle-mobile/paddle-mobile/Executor.swift +++ b/metal/paddle-mobile/paddle-mobile/Executor.swift @@ -14,6 +14,8 @@ import Foundation +let testTo = 41 + public class ResultHolder { public let dim: [Int] public let resultArr: [P] @@ -60,7 +62,7 @@ public class Executor { queue = inQueue for block in inProgram.programDesc.blocks { //block.ops.count - for i in 0..<91 { + for i in 0..<(testTo + 1) { let op = block.ops[i] do { let op = try OpCreator

.shared.creat(device: inDevice, opDesc: op, scope: inProgram.scope) @@ -128,7 +130,7 @@ public class Executor { } // return - self.ops[90].delogOutput() + self.ops[testTo].delogOutput() // self.ops[91].delogOutput() // self.ops[92].delogOutput() // self.ops[93].delogOutput() diff --git a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift index c78055896b4573dbbb715be0dd664b9a3e0a949f..7c1532adf9fbe75ba842ab12fb74d73f90f0899c 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift @@ -63,7 +63,10 @@ class ConcatOp: Operator, ConcatParam

>, Run return o } - print(outputArray.strideArray()) +// 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 99d577221b0c3632c83536c652ac688a5b12641f..b4468af42838d028270e5f62c320ddfeedddee98 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal +++ b/metal/paddle-mobile/paddle-mobile/Operators/Kernels/metal/ReshapeKernel.metal @@ -47,25 +47,62 @@ kernel void reshape(texture2d_array inTexture [[texture(0)] 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]); + 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(ixyzn[0], ixyzn[1]), ixyzn[2])[ixyzn[3]]; +// 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)]], + 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 d132d590b4b3131354244d01ab2c26c0daef052c..d637a176bcda9878b542c39ab6704bc0240a7149 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift @@ -55,5 +55,8 @@ class ReshapeOp: Operator, ReshapeParam

>, // 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()) } } diff --git a/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift b/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift index 387353e236146c43a263951792f22c9827a02287..9b2d98ff945bce1ac59f021a1d0afae0635607db 100644 --- a/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift +++ b/metal/paddle-mobile/paddle-mobile/Operators/TransposeOp.swift @@ -60,6 +60,7 @@ class TransposeOp: Operator, TransposeParam } print(outputArray.strideArray()) + // writeToLibrary(fileName: "transpose_ouput", array: outputArray) }