Merge pull request #842 from codeWorm2015/metal

update

metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift

@@ -381,7 +381,7 @@ public extension MTLTexture {
       for w in 0..<dim.w {
         for sliceIndex in 0..<arrayLength {
           if sliceIndex * 4 + 4 > 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)
             }

metal/paddle-mobile/paddle-mobile/Executor.swift

@@ -14,7 +14,7 @@
 
 import Foundation
 
-let testTo = 48
+let testTo = 93
 
 public class ResultHolder<P: PrecisionType> {
   public let dim: [Int]

metal/paddle-mobile/paddle-mobile/Operators/ConcatOp.swift

@@ -59,15 +59,21 @@ class ConcatOp<P: PrecisionType>: Operator<ConcatKernel<P>, ConcatParam<P>>, 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])
     

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<float, access::read> inTexture [[texture(0)]],
+//                    texture2d_array<float, access::write> 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<float, access::read> inTexture [[texture(0)]],
                     texture2d_array<float, access::write> outTexture [[texture(1)]],
                     constant ReshapeParam &rp [[buffer(0)]],
@@ -32,7 +78,7 @@ kernel void reshape(texture2d_array<float, access::read> 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<float, access::read> 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<float, access::read> inTexture [[texture(0)]
   outTexture.write(r, gid.xy, gid.z);
 }
 
-
-/*
- 
- kernel void reshape(texture2d_array<float, access::read> inTexture [[texture(0)]],
- texture2d_array<float, access::write> 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<half, access::read> inTexture [[texture(0)]],
 //                         texture2d_array<half, access::write> outTexture [[texture(1)]],

metal/paddle-mobile/paddle-mobile/Operators/ReshapeOp.swift

@@ -53,10 +53,11 @@ class ReshapeOp<P: PrecisionType>: Operator<ReshapeKernel<P>, ReshapeParam<P>>,
   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())
+
   }
 }