hollance / forge Goto Github PK

tensorflow::Status run_status = tf_session->Run(
        {{input_layer_name, image_tensor}}, {output_layer_name}, {}, &outputs);

conv = MPSCNNConvolution(device: device,
                             convolutionDescriptor: desc,
                             kernelWeights: weights.pointer,
                             biasTerms: biases?.pointer,
                             flags: .none)

  let layer = MPSCNNConvolution(device: device,
                                convolutionDescriptor: desc,
                                kernelWeights: weightsData.pointer,
                                biasTerms: biasData?.pointer,
                                flags: .none)

let buffer4c = [
// R    G     B   A    R    G     B     A
1.0, 0.0, 0.0, 1.0,  1.0, 0.0, 0.0, 1.0, 
// R    G     B   A    R    G     B     A
1.0, 0.0, 0.0, 1.0,  1.0, 0.0, 0.0, 1.0, 
]

inputImg  = MPSImage(device: device,
                                             numberOfImages: 1,
                                             width: 2,
                                             height: 2,
                                             featureChannels: 3,
                                             array: &buffer4c,
                                             count: 2*2*4)

nums = [1.0,1.0,1.0, 1.0,1.0,1.0,
                1.0,1.0,1.0, 1.0,1.0,1.0]

class Conv2d : NeuralNetwork {
    typealias PredictionType = Float16
    
    var inputImg: MPSImage!
    var outputImg: MPSImage!
    var oid = MPSImageDescriptor(channelFormat: .float16, width: 1, height: 1, featureChannels: 1)
    var conv2d: MPSCNNConvolution
    
    init(device: MTLDevice, inflightBuffers: Int) {
        weightsLoader   = { name, count in ParameterLoaderBundle(name: name, count: count, suffix: "_W", ext: "txt") }
        outputImg       = MPSImage(device: device, imageDescriptor: oid)
        conv2d          = convolution(device: device, kernel: (2, 2), inChannels: 3, outChannels: 1, activation: nil, name: "conv", useBias: false)
    }
    
    func encode(commandBuffer: MTLCommandBuffer, texture: MTLTexture, inflightIndex: Int) {
        conv2d.encode(commandBuffer: commandBuffer, sourceImage: inputImg, destinationImage: outputImg)
    }
    func fetchResult(inflightIndex: Int) -> NeuralNetworkResult<Float16> {
        let probabilities = outputImg.toFloatArray()
        print(probabilities)
        return NeuralNetworkResult<Float16>()
    }
}

    let relu = MPSCNNNeuronReLU(device: device, a : 0.0)
    let input = Input(width: 256, height: 512, channels:3)
    let mbv1_conv_1 = input
        --> Resize(width: 256, height: 512)
        --> Convolution(kernel: (3, 3), channels: 16, stride: (2, 2), padding: .same, activation: relu, useBias: true, name: "0")
        --> DepthwiseConvolution(kernel: (3, 3), stride: (1, 1), activation: nil, useBias: false, name: "1_d")
        --> PointwiseConvolution(channels: 32, stride: (1, 1), activation: relu, useBias: true, name: "1_p")
        --> DepthwiseConvolution(kernel: (3, 3), stride: (2, 2), activation: nil, useBias: false, name: "2_d")
        --> PointwiseConvolution(channels: 64, stride: (1, 1), activation: relu, useBias: true, name: "2_p")
    
    let mbv1_conv_2 = mbv1_conv_1
        --> DepthwiseConvolution(kernel: (3, 3), stride: (2, 2), activation: nil, useBias: false, name: "3_d")
        --> PointwiseConvolution(channels: 128, stride: (1, 1), activation: relu, useBias: true, name: "3_p")
        --> DepthwiseConvolution(kernel: (3, 3), stride: (1, 1), activation: nil, useBias: false, name: "4_d")
        --> PointwiseConvolution(channels: 128, stride: (1, 1), activation: relu, useBias: true, name: "4_p")
    
    let mbv1_conv_3 = mbv1_conv_2
        --> DepthwiseConvolution(kernel: (3, 3), stride: (1, 1), activation: nil, useBias: false, name: "5_d")
        --> PointwiseConvolution(channels: 256, stride: (1, 1), activation: relu, useBias: true, name: "5_p")
        --> DepthwiseConvolution(kernel: (3, 3), stride: (1, 1), activation: nil, useBias: false, name: "6_d")
        --> PointwiseConvolution(channels: 256, stride: (1, 1), activation: relu, useBias: true, name: "6_p")
    
    let mbv1_maxpool = mbv1_conv_1
        --> MaxPooling(kernel: (2, 2), stride: (2, 2), padding: .valid)
    
    let concat = Concatenate([ mbv1_maxpool, mbv1_conv_2, mbv1_conv_3])
    
    let mbv1_conv4 = concat
        --> DepthwiseConvolution(kernel: (3, 3), stride: (1, 1), activation: nil, useBias: false, name: "7_d")
        --> PointwiseConvolution(channels: 32, stride: (1, 1), activation: relu, useBias: true, name: "7_p")

failed assertion `[MPSTemporaryImage prefetchStorageWithCommandBuffer:imageDescriptorList:] Error: the descriptor must be configured with MTLStorageModePrivate'

return MPSImageDescriptor(channelFormat: .float16, width: width,
                              height: height, featureChannels: channels) 

return MPSImageDescriptor(channelFormat: .float16, width: width,
                              height: height, featureChannels: channels,
                              storageMode: .private) // and MTLStorageMode.private instead of .private

dyld: Library not loaded: @rpath/Forge.framework/Forge
  Referenced from: /var/containers/Bundle/Application/xxxxxxxxxxxxxxx/Inception.app/Inception
  Reason: no suitable image found.  Did find:
	/private/var/containers/Bundle/Application/xxxxxxxxxxxxxxx/Inception.app/Frameworks/Forge.framework/Forge: code signing blocked mmap() of '/private/var/containers/Bundle/Application/xxxxxxxxxxxxxx/Inception.app/Frameworks/Forge.framework/Forge'
(lldb)

CodeSign /Users/adamszendrei/Library/Developer/Xcode/DerivedData/Forge-gikysnjirgrtmefxffiekblqjkzy/Build/Products/Debug-iphoneos/Inception.app/Frameworks/Forge.framework
    cd /Users/adamszendrei/ObjDetect/Forge/Examples/Inception
    export CODESIGN_ALLOCATE=/Applications/Xcode.app/Contents/Developer/Toolchains/XcodeDefault.xctoolchain/usr/bin/codesign_allocate
    export PATH="/Applications/Xcode.app/Contents/Developer/Platforms/iPhoneOS.platform/Developer/usr/bin:/Applications/Xcode.app/Contents/Developer/usr/bin:/usr/local/bin:/usr/bin:/bin:/usr/sbin:/sbin"
    
Signing Identity:     "iPhone Developer: Adam Szendrei (762xxxxxxxx)"

    /usr/bin/codesign --force --sign 6xxxxxx --preserve-metadata=identifier,entitlements,flags --timestamp=none /Users/adamszendrei/Library/Developer/Xcode/DerivedData/Forge-gikysnjirgrtmefxffiekblqjkzy/Build/Products/Debug-iphoneos/Inception.app/Frameworks/Forge.framework

Warning: unable to build chain to self-signed root for signer "iPhone Developer: Adam Szendrei (762xxxxxxxx)"

func session(_ session: ARSession, didUpdate frame: ARFrame) {
        let seekingCM = CMTimeMakeWithSeconds(frame.timestamp, 1000000);
        let timestamp = seekingCM
        let deltaTime = timestamp - lastTimestamp
        if fps == -1 || deltaTime >= CMTimeMake(1, Int32(fps)) {
            lastTimestamp = timestamp
            
            if let texture = convertToMTLTexture(pixelBuffer:frame.capturedImage){
                predict(texture: texture)
            }
            
        }
    }

func convertToMTLTexture(pixelBuffer: CVPixelBuffer?) -> MTLTexture? {
        if let textureCache = textureCache,
            let pixelBuffer = pixelBuffer{

            let width = CVPixelBufferGetWidth(pixelBuffer)
            let height = CVPixelBufferGetHeight(pixelBuffer)
            
            var texture: CVMetalTexture?
            CVMetalTextureCacheCreateTextureFromImage(kCFAllocatorDefault, textureCache,
                                                      pixelBuffer, nil, .bgra8Unorm, width, height, 0, &texture)
            if let texture = texture {
                return CVMetalTextureGetTexture(texture)
            }
        }
        return nil
    }

private func show(predictions: [YOLO.Prediction]) {
        DEBUGLOG(message: predictions.count)

        for i in 0..<boundingBoxes.count {
            if i < predictions.count {
                let prediction = predictions[i]
                
                // The predicted bounding box is in the coordinate space of the input
                // image, which is a square image of 416x416 pixels. We want to show it
                // on the video preview, which is as wide as the screen and has a 4:3
                // aspect ratio. The video preview also may be letterboxed at the top
                // and bottom.
                let width = view.bounds.width
                let height = width * 16 / 9
                let scaleX = width / CGFloat(YOLO.inputWidth)
                let scaleY = height / CGFloat(YOLO.inputHeight)
//                let top = (view.bounds.height - height) / 2
                
                // Translate and scale the rectangle to our own coordinate system.
                var rect = prediction.rect
                rect.origin.x *= scaleX
                rect.origin.y *= scaleY
//                rect.origin.y += top
                rect.size.width *= scaleX
                rect.size.height *= scaleY
                
                // Show the bounding box.
                let label = String(format: "%@ %.1f", labels[prediction.classIndex], prediction.score * 100)
                let color = colors[prediction.classIndex]
                boundingBoxes[i].show(frame: rect, label: label, color: color)
                
            } else {
                boundingBoxes[i].hide()
            }
        }
    }

NSDate *start2 = [NSDate date];
    MPSImageDescriptor *desc = [MPSImageDescriptor imageDescriptorWithChannelFormat:MPSImageFeatureChannelFormatFloat32 width:224 height:224 featureChannels:3];
    MPSImage *srcImage = [[MPSImage alloc] initWithDevice:self.device imageDescriptor:desc];
    
    MPSImageDescriptor *desc2 = [MPSImageDescriptor imageDescriptorWithChannelFormat:MPSImageFeatureChannelFormatFloat32 width:224 height:224 featureChannels:64];
    MPSImage *outImage = [[MPSImage alloc] initWithDevice:self.device imageDescriptor:desc2];
    
    id<MTLCommandBuffer> commandBuffer = [self.commandQueue commandBuffer];

    int co = 4*224*224;
    int kernel_size = 3;
    int pad = 1;
    int stride = 1;
    int count = 64*224*224;
    
    float *buf = new float[co];
    for(int i =0;i<co;i++){
        buf[i] = 1.0;
    }
    
    int weight_count = 3*64*kernel_size*kernel_size;
    float *weight = new float[weight_count];
    for(int i =0;i<weight_count;i++){
        weight[i] = 0.123;
    }

    float *bias = new float[64];
    for(int i =0;i<64;i++){
        bias[i] = 1.23456789;
    }
    MTLRegion region = MTLRegionMake3D(0, 0, 0,224,224,1);
    [srcImage.texture replaceRegion:region mipmapLevel:0 slice:0 withBytes:buf bytesPerRow:srcImage.width*4*sizeof(float) bytesPerImage:0];
    
    MPSCNNConvolutionDescriptor *convdesc = [MPSCNNConvolutionDescriptor cnnConvolutionDescriptorWithKernelWidth:kernel_size kernelHeight:kernel_size inputFeatureChannels:3 outputFeatureChannels:64 neuronFilter:nil];
    convdesc.strideInPixelsX = stride;
    convdesc.strideInPixelsY = stride;
    convdesc.groups = 1;
    
    MPSCNNConvolution *conv = [[MPSCNNConvolution alloc] initWithDevice:self.device convolutionDescriptor:convdesc kernelWeights:weight biasTerms:bias flags:MPSCNNConvolutionFlagsNone];
    MPSOffset offset;
    offset.x = 0;
    offset.y = 0;
    offset.z = 0;
    conv.offset = offset;
    
    
    [conv encodeToCommandBuffer:commandBuffer sourceImage:srcImage destinationImage:outImage];
    NSTimeInterval localtime2 = [[NSDate date] timeIntervalSinceDate:start2] * 1000;
    cout << "data init used " << localtime2 << "ms" << endl;
    
    
    NSDate *start = [NSDate date];

    [commandBuffer commit];
    [commandBuffer waitUntilCompleted];
    
    delete [] buf;
    delete [] weight;
    delete [] bias;
    NSTimeInterval localtime = [[NSDate date] timeIntervalSinceDate:start] * 1000;

    cout << "gpu calc used " << localtime << "ms" << endl;

id <MTLComputePipelineState> pipline = self.pipelineShaderTex;
    
    int co = 4*224*224;
    int kernel_size = 3;
    int pad = 1;
    int stride = 1;
    int count = 64*224*224;
    
    float *buf = new float[co];
    for(int i =0;i<co;i++){
        buf[i] = 1.0;
    }
    
    int weight_count = 4*64*kernel_size*kernel_size;
    float *weight = new float[weight_count];
    for(int i =0;i<weight_count;i++){
        weight[i] = i%4 == 3 ? 0 : 0.123;
    }

    float *bias = new float[64];
    for(int i =0;i<64;i++){
        bias[i] = 1.23456789;
    }
    
    MetalConvolutionParameter param;
    param.count = count;
    param.padSize = pad;
    param.kernelSize = kernel_size;
    param.stride = stride;
    param.inputChannel = 3;
    param.outputChannel = 64;
    param.inputWidth = 224;
    param.inputHeight = 224;
    param.outputWidth = 224;
    param.outputHeight = 224;
    
    MTLTextureDescriptor *indesc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA32Float width:224 height:224 mipmapped:NO];
    indesc.textureType = MTLTextureType2D;
    
    MTLTextureDescriptor *outdesc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA32Float width:224 height:224 mipmapped:NO];
    outdesc.textureType = MTLTextureType2DArray;
    outdesc.arrayLength = 64/4;
    
    MTLTextureDescriptor *weightdesc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA32Float width:3 height:3 mipmapped:NO];
    weightdesc.textureType = MTLTextureType2DArray;
    weightdesc.arrayLength = 64;

    MTLTextureDescriptor *biasdesc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA32Float width:1 height:1 mipmapped:NO];
    biasdesc.textureType = MTLTextureType2DArray;
    biasdesc.arrayLength = 64/4;
    
    if(!self.inTexture){
        self.inTexture = [self.device newTextureWithDescriptor:indesc];
        self.outTexture = [self.device newTextureWithDescriptor:outdesc];
        self.weightTexture = [self.device newTextureWithDescriptor:weightdesc];
        self.biasTexture = [self.device newTextureWithDescriptor:biasdesc];
        
        [self.inTexture replaceRegion:MTLRegionMake3D(0, 0, 0, 224, 224, 1) mipmapLevel:0 slice:0 withBytes:buf bytesPerRow:224*4*sizeof(float) bytesPerImage:0];
        for(int i =0;i<weightdesc.arrayLength;i++){
            [self.weightTexture replaceRegion:MTLRegionMake3D(0, 0, 0, 3, 3, 1) mipmapLevel:0 slice:i withBytes:weight+3*3*4*i bytesPerRow:3*4*sizeof(float) bytesPerImage:0];
            
        }
        for(int i =0;i<biasdesc.arrayLength;i++){
            [self.biasTexture replaceRegion:MTLRegionMake3D(0, 0, 0, 1, 1, 1) mipmapLevel:0 slice:i withBytes:bias+4*i bytesPerRow:1*4*sizeof(float) bytesPerImage:0];
        }
    }
    id<MTLBuffer> parambuffer = [self.device newBufferWithBytes:&param length:sizeof(param) options:MTLResourceCPUCacheModeDefaultCache];

    id<MTLCommandBuffer> commandBuffer = [self.commandQueue commandBuffer];
    id<MTLComputeCommandEncoder> encoder = [commandBuffer computeCommandEncoder];
    [encoder setComputePipelineState:pipline];
    [encoder setTexture:self.inTexture atIndex:0];
    [encoder setTexture:self.outTexture atIndex:1];
    [encoder setTexture:self.weightTexture atIndex:2];
    [encoder setTexture:self.biasTexture atIndex:3];
    [encoder setBuffer:parambuffer offset:0 atIndex:0];
    
    MTLSize threadsPerGroups = MTLSizeMake(32, 16, 1);
    MTLSize threadGroups = MTLSizeMake((224 + threadsPerGroups.width -1 ) / threadsPerGroups.width,
                                       (224 + threadsPerGroups.height -1 ) / threadsPerGroups.height, 16);
    
    [encoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadsPerGroups];
    [encoder endEncoding];
    
    NSDate *start = [NSDate date];

    [commandBuffer commit];
    [commandBuffer waitUntilCompleted];
    
    delete [] buf;
    delete [] weight;
    delete [] bias;
    NSTimeInterval localtime = [[NSDate date] timeIntervalSinceDate:start] * 1000;
    cout << "Time used " << localtime << "ms" << endl;

kernel void convolutionForwardTexture(texture2d<float, access::read> inTexture [[texture(0)]],
                                      texture2d_array<float, access::write> outTexture [[texture(1)]],
                                      texture2d_array<float, access::read> weights [[ texture(2) ]],
                                      texture2d_array<float, access::read> bias [[ texture(3) ]],
                                      const device MetalConvolutionParameter *convolvParams [[ buffer(0) ]],
                                      ushort3 gid [[ thread_position_in_grid ]]){
    if(gid.x>=224||gid.y>=224){
        return;
    }
    
    float total = 0;
    float total2 = 0;
    float total3 = 0;
    float total4 = 0;
    
    float4 k,input;
    int slice = gid.z;
    for(int kh =0;kh<3;kh++){
        for(int kw =0;kw<3;kw++) {
            k = weights.read(uint2(kw,kh),slice*4);
            input = inTexture.read(uint2(0,0));
            total+=dot(k,input);
            
            k = weights.read(uint2(kw,kh),slice*4+1);
            input = inTexture.read(uint2(0,0));
            total2+=dot(k,input);
            
            k = weights.read(uint2(kw,kh),slice*4+2);
            input = inTexture.read(uint2(0,0));
            total3+=dot(k,input);
            
            k = weights.read(uint2(kw,kh),slice*4+3);
            input = inTexture.read(uint2(0,0));
            total4+=dot(k,input);
        }
    }
    
    float4 output = float4(total,total2,total3,total4) + bias.read(uint2(0,0),slice);
    outTexture.write(output,uint2(gid.x,gid.y),gid.z);
    
}