Detection good, split up col images

Joseph Redmon

2014-12-04 1edcf73a73d2007afc61289245763f5cf0c29e10

Detection good, split up col images

 src/axpy.cl

@@ -13,7 +13,7 @@
__kernel void mask(int n, __global float *x, __global float *mask, int mod)
{
    int i = get_global_id(0);
    x[i] = (mask[(i/mod)*mod]) ? x[i] : 0;
    x[i] = (mask[(i/mod)*mod] || i%mod == 0) ? x[i] : 0;
}

__kernel void copy(int N, __global float *X, int OFFX, int INCX, __global float *Y, int OFFY, int INCY)

 src/cnn.c

@@ -36,6 +36,7 @@

void test_convolutional_layer()
{
/*
    int i;
    image dog = load_image("data/dog.jpg",224,224);
    network net = parse_network_cfg("cfg/convolutional.cfg");
@@ -72,6 +73,7 @@

    float *gpu_del = calloc(del_size, sizeof(float));
    memcpy(gpu_del, get_network_delta_layer(net, 0), del_size*sizeof(float));
    */

    /*
       start = clock();
@@ -97,6 +99,7 @@
     */
}

/*
void test_col2im()
{
    float col[] =  {1,2,1,2,
@@ -116,13 +119,12 @@
    int ksize = 3;
    int stride = 1;
    int pad = 0;
    col2im_gpu(col, batch,
            channels,  height,  width,
            ksize,  stride, pad, im);
    //col2im_gpu(col, batch,
    //        channels,  height,  width,
    //        ksize,  stride, pad, im);
    int i;
    for(i = 0; i < 16; ++i)printf("%f,", im[i]);
    printf("\n");
    /*
       float data_im[] = {
       1,2,3,4,
       5,6,7,8,
@@ -134,8 +136,8 @@
       ksize,   stride,  pad, data_col) ;
       for(i = 0; i < 18; ++i)printf("%f,", data_col[i]);
       printf("\n");
     */
}
*/

#endif

@@ -158,7 +160,7 @@
    int i;
    clock_t start = clock(), end;
    for(i = 0; i < 1000; ++i){
        im2col_cpu(dog.data,1, dog.c,  dog.h,  dog.w,  size,  stride, 0, matrix);
        //im2col_cpu(dog.data,1, dog.c,  dog.h,  dog.w,  size,  stride, 0, matrix);
        gemm(0,0,n,mw,mh,1,filters,mh,matrix,mw,1,edge.data,mw);
    }
    end = clock();
@@ -175,6 +177,7 @@

void verify_convolutional_layer()
{
/*
    srand(0);
    int i;
    int n = 1;
@@ -225,6 +228,7 @@
    printf("%f %f\n", avg_image_layer(mj1,0), avg_image_layer(mj2,0));
    show_image(mj1, "forward jacobian");
    show_image(mj2, "backward jacobian");
    */
}

void test_load()
@@ -446,7 +450,7 @@
        for(c = 0; c < 8; ++c){
            j = (r*8 + c) * 5;
            printf("Prob: %f\n", box[j]);
            if(box[j] > .999){
            if(box[j] > .05){
                int d = 256/8;
                int y = r*d+box[j+1]*d;
                int x = c*d+box[j+2]*d;
@@ -465,7 +469,7 @@

void test_detection()
{
    network net = parse_network_cfg("cfg/detnet_test.cfg");
    network net = parse_network_cfg("cfg/detnet.test");
    srand(2222222);
    clock_t time;
    char filename[256];
@@ -726,7 +730,7 @@
    float *matrix = calloc(msize, sizeof(float));
    int i;
    for(i = 0; i < 1000; ++i){
        im2col_cpu(test.data,1,  c,  h,  w,  size,  stride, 0, matrix);
        //im2col_cpu(test.data,1,  c,  h,  w,  size,  stride, 0, matrix);
        //image render = float_to_image(mh, mw, mc, matrix);
    }
}
@@ -782,13 +786,59 @@
#endif
}

void test_correct_alexnet()
{
    char **labels = get_labels("/home/pjreddie/data/imagenet/cls.labels.list");
    list *plist = get_paths("/data/imagenet/cls.train.list");
    char **paths = (char **)list_to_array(plist);
    printf("%d\n", plist->size);
    clock_t time;
    int count = 0;

    srand(222222);
    network net = parse_network_cfg("cfg/alexnet.test");
    printf("Learning Rate: %g, Momentum: %g, Decay: %g\n", net.learning_rate, net.momentum, net.decay);
    int imgs = 1000/net.batch+1;
    imgs = 1;

    while(++count <= 5){
        time=clock();
        data train = load_data_random(imgs*net.batch, paths, plist->size, labels, 1000, 256, 256);
        //translate_data_rows(train, -144);
        normalize_data_rows(train);
        printf("Loaded: %lf seconds\n", sec(clock()-time));
        time=clock();
        float loss = train_network_data_cpu(net, train, imgs);
        printf("%d: %f, %lf seconds, %d images\n", count, loss, sec(clock()-time), imgs*net.batch);
        free_data(train);
    }
#ifdef GPU
    count = 0;
    srand(222222);
    net = parse_network_cfg("cfg/alexnet.test");
    while(++count <= 5){
        time=clock();
        data train = load_data_random(imgs*net.batch, paths, plist->size, labels, 1000, 256, 256);
        //translate_data_rows(train, -144);
        normalize_data_rows(train);
        printf("Loaded: %lf seconds\n", sec(clock()-time));
        time=clock();
        float loss = train_network_data_gpu(net, train, imgs);
        printf("%d: %f, %lf seconds, %d images\n", count, loss, sec(clock()-time), imgs*net.batch);
        free_data(train);
    }
#endif
}

void test_server()
{
    server_update();
    network net = parse_network_cfg("cfg/alexnet.test");
    server_update(net);
}
void test_client()
{
    client_update();
    network net = parse_network_cfg("cfg/alexnet.test");
    client_update(net);
}

int main(int argc, char *argv[])
@@ -801,7 +851,7 @@
    else if(0==strcmp(argv[1], "detection")) train_detection_net();
    else if(0==strcmp(argv[1], "asirra")) train_asirra();
    else if(0==strcmp(argv[1], "nist")) train_nist();
    else if(0==strcmp(argv[1], "test_correct")) test_gpu_net();
    else if(0==strcmp(argv[1], "test_correct")) test_correct_alexnet();
    else if(0==strcmp(argv[1], "test")) test_imagenet();
    else if(0==strcmp(argv[1], "server")) test_server();
    else if(0==strcmp(argv[1], "client")) test_client();

 src/col2im.c

@@ -1,21 +1,21 @@
#include <stdio.h>
#include <math.h>
inline void col2im_add_pixel(float *im, int height, int width, int channels,
                        int b, int row, int col, int channel, int pad, float val)
                        int row, int col, int channel, int pad, float val)
{
    row -= pad;
    col -= pad;

    if (row < 0 || col < 0 ||
        row >= height || col >= width) return;
    im[col + width*(row + height*(channel+b*channels))] += val;
    im[col + width*(row + height*channel)] += val;
}
//This one might be too, can't remember.
void col2im_cpu(float* data_col, int batch,
void col2im_cpu(float* data_col,
         int channels,  int height,  int width,
         int ksize,  int stride, int pad, float* data_im) 
{
    int b,c,h,w;
    int c,h,w;
    int height_col = (height - ksize) / stride + 1;
    int width_col = (width - ksize) / stride + 1;
    if (pad){
@@ -24,8 +24,6 @@
        pad = ksize/2;
    }
    int channels_col = channels * ksize * ksize;
    int col_size = height_col*width_col*channels_col;
    for(b = 0; b < batch; ++b){
        for (c = 0; c < channels_col; ++c) {
            int w_offset = c % ksize;
            int h_offset = (c / ksize) % ksize;
@@ -34,11 +32,10 @@
                for (w = 0; w < width_col; ++w) {
                    int im_row = h_offset + h * stride;
                    int im_col = w_offset + w * stride;
                    int col_index = (c * height_col + h) * width_col + w + b*col_size;
                int col_index = (c * height_col + h) * width_col + w;
                    double val = data_col[col_index];
                    col2im_add_pixel(data_im, height, width, channels,
                            b, im_row, im_col, c_im, pad, val);
                }
                        im_row, im_col, c_im, pad, val);
            }
        }
    }
@@ -60,7 +57,7 @@
    return im2col_kernel;
}

void col2im_ongpu(cl_mem data_col,  int batch,
void col2im_ongpu(cl_mem data_col,  int offset,
         int channels,  int height,  int width,
         int ksize,  int stride,  int pad, cl_mem data_im)
{
@@ -70,7 +67,7 @@

    cl_uint i = 0;
    cl.error = clSetKernelArg(kernel, i++, sizeof(data_col), (void*) &data_col);
    cl.error = clSetKernelArg(kernel, i++, sizeof(batch), (void*) &batch);
    cl.error = clSetKernelArg(kernel, i++, sizeof(offset), (void*) &offset);
    cl.error = clSetKernelArg(kernel, i++, sizeof(channels), (void*) &channels);
    cl.error = clSetKernelArg(kernel, i++, sizeof(height), (void*) &height);
    cl.error = clSetKernelArg(kernel, i++, sizeof(width), (void*) &width);
@@ -80,13 +77,14 @@
    cl.error = clSetKernelArg(kernel, i++, sizeof(data_im), (void*) &data_im);
    check_error(cl);

    size_t global_size = channels*height*width*batch;
    size_t global_size = channels*height*width;

    cl.error = clEnqueueNDRangeKernel(queue, kernel, 1, 0,
            &global_size, 0, 0, 0, 0);
    check_error(cl);
}

/*
void col2im_gpu(float *data_col,  int batch,
         int channels,  int height,  int width,
         int ksize,  int stride,  int pad, float *data_im) 
@@ -107,5 +105,6 @@
    clReleaseMemObject(col_gpu);
    clReleaseMemObject(im_gpu);
}
 */

#endif

 src/col2im.cl

@@ -1,4 +1,4 @@
__kernel void col2im(__global float *data_col, int batch,
__kernel void col2im(__global float *data_col, int offset,
        int channels, int height, int width,
        int ksize, int stride, int pad, __global float *data_im)
{
@@ -18,33 +18,26 @@
    int h = id%height + pad;
    id /= height;
    int c = id%channels;
    id /= channels;
    int b = id%batch;

    //int w_start = (w<ksize)?0:(w-ksize)/stride + 1;
    int w_start = (w-ksize+stride)/stride;
    int w_end = w/stride + 1;
    //w_end = (width_col < w_end) ? width_col : w_end;

    int h_start = (h-ksize+stride)/stride;
    //int h_start = (h-ksize)/stride+1;
    int h_end = h/stride + 1;
    //h_end = (height_col < h_end) ? height_col : h_end;

    int rows = channels * ksize * ksize;
    int cols = height_col*width_col;
    int offset = (c*ksize*ksize + h * ksize + w)*height_col*width_col;
    offset += b*cols*rows;
    int col_offset = (c*ksize*ksize + h * ksize + w)*height_col*width_col;
    int h_coeff = (1-stride*ksize*height_col)*width_col;
    int w_coeff = 1-stride*height_col*width_col;
    float val = 0;
    int h_col, w_col;
    for(h_col = h_start; h_col < h_end; ++h_col){
        for(w_col = w_start; w_col < w_end; ++w_col){
            int col_index = offset +h_col*h_coeff + w_col*w_coeff;
            int col_index = col_offset +h_col*h_coeff + w_col*w_coeff;
            float part = (w_col < 0 || h_col < 0 || h_col >= height_col || w_col >= width_col) ? 0 : data_col[col_index];
            val += part;
        }
    }
    data_im[index] = val;
    data_im[index+offset] = val;
}

 src/convolutional_layer.c

@@ -65,7 +65,7 @@
    layer->bias_updates = calloc(n, sizeof(float));
    layer->bias_momentum = calloc(n, sizeof(float));
    float scale = 1./(size*size*c);
    scale = .05;
    scale = .01;
    for(i = 0; i < c*n*size*size; ++i) layer->filters[i] = scale*2*(rand_uniform()-.5);
    for(i = 0; i < n; ++i){
        //layer->biases[i] = rand_normal()*scale + scale;
@@ -74,7 +74,7 @@
    int out_h = convolutional_out_height(*layer);
    int out_w = convolutional_out_width(*layer);

    layer->col_image = calloc(layer->batch*out_h*out_w*size*size*c, sizeof(float));
    layer->col_image = calloc(out_h*out_w*size*size*c, sizeof(float));
    layer->output = calloc(layer->batch*out_h * out_w * n, sizeof(float));
    layer->delta  = calloc(layer->batch*out_h * out_w * n, sizeof(float));
    #ifdef GPU
@@ -86,7 +86,7 @@
    layer->bias_updates_cl = cl_make_array(layer->bias_updates, n);
    layer->bias_momentum_cl = cl_make_array(layer->bias_momentum, n);

    layer->col_image_cl = cl_make_array(layer->col_image, layer->batch*out_h*out_w*size*size*c);
    layer->col_image_cl = cl_make_array(layer->col_image, out_h*out_w*size*size*c);
    layer->delta_cl = cl_make_array(layer->delta, layer->batch*out_h*out_w*n);
    layer->output_cl = cl_make_array(layer->output, layer->batch*out_h*out_w*n);
    #endif
@@ -106,7 +106,7 @@
    int out_w = convolutional_out_width(*layer);

    layer->col_image = realloc(layer->col_image,
                                layer->batch*out_h*out_w*layer->size*layer->size*layer->c*sizeof(float));
                                out_h*out_w*layer->size*layer->size*layer->c*sizeof(float));
    layer->output = realloc(layer->output,
                                layer->batch*out_h * out_w * layer->n*sizeof(float));
    layer->delta  = realloc(layer->delta,
@@ -143,13 +143,13 @@
    float *b = layer.col_image;
    float *c = layer.output;

    im2col_cpu(in, layer.batch, layer.c, layer.h, layer.w, 
        layer.size, layer.stride, layer.pad, b);

    for(i = 0; i < layer.batch; ++i){
        im2col_cpu(in, layer.c, layer.h, layer.w, 
            layer.size, layer.stride, layer.pad, b);
        gemm(0,0,m,n,k,1,a,k,b,n,1,c,n);
        b += k*n;
        c += n*m;
        in += layer.c*layer.h*layer.w;
    }
    activate_array(layer.output, m*n*layer.batch, layer.activation);
}
@@ -166,7 +166,7 @@
    }
}

void backward_convolutional_layer(convolutional_layer layer, float *delta)
void backward_convolutional_layer(convolutional_layer layer, float *in, float *delta)
{
    int i;
    int m = layer.n;
@@ -176,35 +176,28 @@
    gradient_array(layer.output, m*k*layer.batch, layer.activation, layer.delta);
    learn_bias_convolutional_layer(layer);

    float *a = layer.delta;
    if(delta) memset(delta, 0, layer.batch*layer.h*layer.w*layer.c*sizeof(float));

    for(i = 0; i < layer.batch; ++i){
        float *a = layer.delta + i*m*k;
    float *b = layer.col_image;
    float *c = layer.filter_updates;

    for(i = 0; i < layer.batch; ++i){
        float *im = in+i*layer.c*layer.h*layer.w;

        im2col_cpu(im, layer.c, layer.h, layer.w, 
                layer.size, layer.stride, layer.pad, b);
        gemm(0,1,m,n,k,1,a,k,b,k,1,c,n);
        a += m*k;
        b += k*n;
    }

    if(delta){
        m = layer.size*layer.size*layer.c;
        k = layer.n;
        n = convolutional_out_height(layer)*
            convolutional_out_width(layer);

        a = layer.filters;
        b = layer.delta;
            b = layer.delta + i*m*k;
        c = layer.col_image;

        for(i = 0; i < layer.batch; ++i){
            gemm(1,0,m,n,k,1,a,m,b,n,0,c,n);
            b += k*n;
            c += m*n;
            gemm(1,0,n,k,m,1,a,n,b,k,0,c,k);

            col2im_cpu(layer.col_image, layer.c,  layer.h,  layer.w,  layer.size,  layer.stride, layer.pad, delta+i*layer.c*layer.h*layer.w);
        }

        memset(delta, 0, layer.batch*layer.h*layer.w*layer.c*sizeof(float));

        col2im_cpu(layer.col_image, layer.batch, layer.c,  layer.h,  layer.w,  layer.size,  layer.stride, layer.pad, delta);
    }
}

@@ -354,36 +347,17 @@

    bias_output_gpu(layer);

    #ifdef TIMEIT
    clock_t time = clock();
    printf("Forward\n");
    #endif

    im2col_ongpu(in, layer.batch, layer.c,  layer.h,  layer.w,  layer.size,  layer.stride, layer.pad, layer.col_image_cl);

    #ifdef TIMEIT
    clFinish(cl.queue);
    printf("Im2col %f\n", sec(clock()-time));
    time = clock();
    #endif

    for(i = 0; i < layer.batch; ++i){
        im2col_ongpu(in, i*layer.c*layer.h*layer.w, layer.c,  layer.h,  layer.w,  layer.size,  layer.stride, layer.pad, layer.col_image_cl);
        cl_mem a = layer.filters_cl;
        cl_mem b = layer.col_image_cl;
        cl_mem c = layer.output_cl;
        gemm_ongpu_offset(0,0,m,n,k,1.,a,0,k,b,i*k*n,n,1.,c,i*m*n,n);
        gemm_ongpu_offset(0,0,m,n,k,1.,a,0,k,b,0,n,1.,c,i*m*n,n);
    }
    #ifdef TIMEIT
    clFinish(cl.queue);
    printf("Gemm %f\n", sec(clock()-time));
    #endif
    activate_array_ongpu(layer.output_cl, m*n*layer.batch, layer.activation);
    #ifdef TIMEIT
    cl_read_array(layer.output_cl, layer.output, m*n*layer.batch);
    #endif
}

void backward_convolutional_layer_gpu(convolutional_layer layer, cl_mem delta_cl)
void backward_convolutional_layer_gpu(convolutional_layer layer, cl_mem in, cl_mem delta_cl)
{
    int i;
    int m = layer.n;
@@ -393,30 +367,26 @@
    gradient_array_ongpu(layer.output_cl, m*k*layer.batch, layer.activation, layer.delta_cl);
    learn_bias_convolutional_layer_ongpu(layer);

    if(delta_cl) scal_ongpu(layer.batch*layer.h*layer.w*layer.c, 0, delta_cl, 1);

    for(i = 0; i < layer.batch; ++i){
        cl_mem a = layer.delta_cl;
        cl_mem b = layer.col_image_cl;
        cl_mem c = layer.filter_updates_cl;

        gemm_ongpu_offset(0,1,m,n,k,1,a,i*m*k,k,b,i*k*n,k,1,c,0,n);
    }
        im2col_ongpu(in, i*layer.c*layer.h*layer.w, layer.c,  layer.h,  layer.w,  layer.size,  layer.stride, layer.pad, layer.col_image_cl);
        gemm_ongpu_offset(0,1,m,n,k,1,a,i*m*k,k,b,0,k,1,c,0,n);

    if(delta_cl){
        m = layer.size*layer.size*layer.c;
        k = layer.n;
        n = convolutional_out_height(layer)*
            convolutional_out_width(layer);

        for(i = 0; i < layer.batch; ++i){
            cl_mem a = layer.filters_cl;
            cl_mem b = layer.delta_cl;
            cl_mem c = layer.col_image_cl;

            gemm_ongpu_offset(1,0,m,n,k,1,a,0,m,b,i*k*n,n,0,c,i*m*n,n);
        }
            gemm_ongpu_offset(1,0,n,k,m,1,a,0,n,b,i*k*m,k,0,c,0,k);

        scal_ongpu(layer.batch*layer.h*layer.w*layer.c,0,delta_cl, 1);
        col2im_ongpu(layer.col_image_cl, layer.batch, layer.c,  layer.h,  layer.w,  layer.size,  layer.stride, layer.pad, delta_cl);
            col2im_ongpu(layer.col_image_cl, i*layer.c*layer.h*layer.w, layer.c,  layer.h,  layer.w,  layer.size,  layer.stride, layer.pad, delta_cl);
        }
    }
}


 src/convolutional_layer.h

@@ -47,7 +47,7 @@

#ifdef GPU
void forward_convolutional_layer_gpu(convolutional_layer layer, cl_mem in);
void backward_convolutional_layer_gpu(convolutional_layer layer, cl_mem delta_cl);
void backward_convolutional_layer_gpu(convolutional_layer layer, cl_mem in, cl_mem delta_cl);
void update_convolutional_layer_gpu(convolutional_layer layer);
void push_convolutional_layer(convolutional_layer layer);
#endif
@@ -58,7 +58,7 @@
void update_convolutional_layer(convolutional_layer layer);
image *visualize_convolutional_layer(convolutional_layer layer, char *window, image *prev_filters);

void backward_convolutional_layer(convolutional_layer layer, float *delta);
void backward_convolutional_layer(convolutional_layer layer, float *in, float *delta);

image get_convolutional_image(convolutional_layer layer);
image get_convolutional_delta(convolutional_layer layer);

 src/cost_layer.c

@@ -52,6 +52,7 @@
        }
    }
    *(layer.output) = dot_cpu(layer.batch*layer.inputs, layer.delta, 1, layer.delta, 1);
    //printf("cost: %f\n", *layer.output);
}

void backward_cost_layer(const cost_layer layer, float *input, float *delta)
@@ -105,7 +106,7 @@

    cl_read_array(layer.delta_cl, layer.delta, layer.batch*layer.inputs);
    *(layer.output) = dot_cpu(layer.batch*layer.inputs, layer.delta, 1, layer.delta, 1);
    //printf("%f\n", *layer.output);
    //printf("cost: %f\n", *layer.output);
}

void backward_cost_layer_gpu(const cost_layer layer, cl_mem input, cl_mem delta)

 src/im2col.c

@@ -1,23 +1,23 @@
#include "mini_blas.h"
#include <stdio.h>
inline float im2col_get_pixel(float *im, int height, int width, int channels,
                        int b, int row, int col, int channel, int pad)
                        int row, int col, int channel, int pad)
{
    row -= pad;
    col -= pad;

    if (row < 0 || col < 0 ||
        row >= height || col >= width) return 0;
    return im[col + width*(row + height*(channel+b*channels))];
    return im[col + width*(row + height*channel)];
}

//From Berkeley Vision's Caffe!
//https://github.com/BVLC/caffe/blob/master/LICENSE
void im2col_cpu(float* data_im,  int batch,
void im2col_cpu(float* data_im,
     int channels,  int height,  int width,
     int ksize,  int stride, int pad, float* data_col) 
{
    int c,h,w,b;
    int c,h,w;
    int height_col = (height - ksize) / stride + 1;
    int width_col = (width - ksize) / stride + 1;
    if (pad){
@@ -26,8 +26,6 @@
        pad = ksize/2;
    }
    int channels_col = channels * ksize * ksize;
    int col_size = height_col*width_col*channels_col;
    for (b = 0; b < batch; ++b) {
        for (c = 0; c < channels_col; ++c) {
            int w_offset = c % ksize;
            int h_offset = (c / ksize) % ksize;
@@ -36,10 +34,9 @@
                for (w = 0; w < width_col; ++w) {
                    int im_row = h_offset + h * stride;
                    int im_col = w_offset + w * stride;
                    int col_index = (c * height_col + h) * width_col + w + b*col_size;
                int col_index = (c * height_col + h) * width_col + w;
                    data_col[col_index] = im2col_get_pixel(data_im, height, width, channels,
                            b, im_row, im_col, c_im, pad);
                }
                        im_row, im_col, c_im, pad);
            }
        }
    }
@@ -74,7 +71,7 @@
}


void im2col_ongpu(cl_mem data_im,  int batch,
void im2col_ongpu(cl_mem data_im, int offset,
         int channels,  int height,  int width,
         int ksize,  int stride,  int pad, cl_mem data_col)
{
@@ -95,7 +92,7 @@

    cl_uint i = 0;
    cl.error = clSetKernelArg(kernel, i++, sizeof(data_im), (void*) &data_im);
    cl.error = clSetKernelArg(kernel, i++, sizeof(batch), (void*) &batch);
    cl.error = clSetKernelArg(kernel, i++, sizeof(offset), (void*) &offset);
    cl.error = clSetKernelArg(kernel, i++, sizeof(channels), (void*) &channels);
    cl.error = clSetKernelArg(kernel, i++, sizeof(height), (void*) &height);
    cl.error = clSetKernelArg(kernel, i++, sizeof(width), (void*) &width);
@@ -104,14 +101,15 @@
    cl.error = clSetKernelArg(kernel, i++, sizeof(data_col), (void*) &data_col);
    check_error(cl);

    size_t global_size = batch*channels_col*height_col*width_col;
    size_t global_size = channels_col*height_col*width_col;

    cl.error = clEnqueueNDRangeKernel(queue, kernel, 1, 0,
            &global_size, 0, 0, 0, 0);
    check_error(cl);
}

void im2col_gpu(float *data_im,  int batch,
/*
   void im2col_gpu(float *data_im,
         int channels,  int height,  int width,
         int ksize,  int stride,  int pad, float *data_col) 
{
@@ -144,5 +142,6 @@
    clReleaseMemObject(col_gpu);
    clReleaseMemObject(im_gpu);
}
 */

#endif

 src/im2col.cl

@@ -1,9 +1,9 @@

__kernel void im2col_pad(__global float *im,  int batch,
__kernel void im2col_pad(__global float *im,  int offset,
     int channels,  int height,  int width,
     int ksize,  int stride, __global float *data_col)
{
    int c,h,w,b;
    int c,h,w;
    int height_col = 1 + (height-1) / stride;
    int width_col = 1 + (width-1) / stride;
    int channels_col = channels * ksize * ksize;
@@ -18,8 +18,6 @@
    id /= height_col;
    c = id % channels_col;
    id /= channels_col;
    b = id % batch;
    id /= batch;

    int col_size = height_col*width_col*channels_col;
    int w_offset = c % ksize;
@@ -28,17 +26,17 @@
    int im_row = h_offset + h * stride - pad;
    int im_col = w_offset + w * stride - pad;

    int im_index = im_col + width*(im_row + height*(im_channel+b*channels));
    int im_index = offset + im_col + width*(im_row + height*im_channel);
    float val = (im_row < 0 || im_col < 0 || im_row >= height || im_col >= width) ? 0 : im[im_index];

    data_col[col_index] = val;
}

__kernel void im2col_nopad(__global float *im,  int batch,
__kernel void im2col_nopad(__global float *im,  int offset,
        int channels,  int height,  int width,
        int ksize,  int stride, __global float *data_col)
{
    int c,h,w,b;
    int c,h,w;
    int height_col = (height - ksize) / stride + 1;
    int width_col = (width - ksize) / stride + 1;
    int channels_col = channels * ksize * ksize;
@@ -51,8 +49,6 @@
    id /= height_col;
    c = id % channels_col;
    id /= channels_col;
    b = id % batch;
    id /= batch;

    int col_size = height_col*width_col*channels_col;
    int w_offset = c % ksize;
@@ -61,7 +57,7 @@
    int im_row = h_offset + h * stride;
    int im_col = w_offset + w * stride;

    int im_index = im_col + width*(im_row + height*(im_channel+b*channels));
    int im_index = offset + im_col + width*(im_row + height*im_channel);
    float val = (im_row < 0 || im_col < 0 || im_row >= height || im_col >= width) ? 0 : im[im_index];

    data_col[col_index] = val;

 src/image.c

@@ -9,14 +9,14 @@
    int i, c;
    for(c = 0; c < a.c; ++c){
        for(i = x1; i < x2; ++i){
            a.data[i + y1*a.w + c*a.w*a.h] = 0;
            a.data[i + y2*a.w + c*a.w*a.h] = 0;
            a.data[i + y1*a.w + c*a.w*a.h] = (c==0)?1:-1;
            a.data[i + y2*a.w + c*a.w*a.h] = (c==0)?1:-1;
        }
    }
    for(c = 0; c < a.c; ++c){
        for(i = y1; i < y2; ++i){
            a.data[x1 + i*a.w + c*a.w*a.h] = 0;
            a.data[x2 + i*a.w + c*a.w*a.h] = 0;
            a.data[x1 + i*a.w + c*a.w*a.h] = (c==0)?1:-1;
            a.data[x2 + i*a.w + c*a.w*a.h] = (c==0)?1:-1;
        }
    }
}

 src/mini_blas.h

@@ -15,18 +15,18 @@
void copy_ongpu(int N, cl_mem X, int INCX, cl_mem Y, int INCY);
void copy_ongpu_offset(int N, cl_mem X, int OFFX, int INCX, cl_mem Y, int OFFY, int INCY);
void scal_ongpu(int N, float ALPHA, cl_mem X, int INCX);
void im2col_ongpu(cl_mem data_im, int batch,
void im2col_ongpu(cl_mem data_im, int offset,
         int channels, int height, int width,
         int ksize, int stride, int pad, cl_mem data_col);

void col2im_gpu(float *data_col,  int batch,
void col2im_gpu(float *data_col,  int offset,
         int channels,  int height,  int width,
         int ksize,  int stride,  int pad, float *data_im);
void col2im_ongpu(cl_mem data_col, int batch,
        int channels, int height, int width,
        int ksize, int stride, int pad, cl_mem data_im);

void im2col_gpu(float *data_im, int batch,
void im2col_gpu(float *data_im,
         int channels, int height, int width,
         int ksize, int stride, int pad, float *data_col);

@@ -43,11 +43,11 @@
        cl_mem C_gpu, int ldc);
#endif

void im2col_cpu(float* data_im, int batch,
void im2col_cpu(float* data_im,
    int channels, int height, int width,
    int ksize, int stride, int pad, float* data_col);

void col2im_cpu(float* data_col, int batch,
void col2im_cpu(float* data_col,
        int channels, int height, int width,
        int ksize, int stride, int pad, float* data_im);


 src/network.c

@@ -213,7 +213,7 @@
        }
        if(net.types[i] == CONVOLUTIONAL){
            convolutional_layer layer = *(convolutional_layer *)net.layers[i];
            backward_convolutional_layer(layer, prev_delta);
            backward_convolutional_layer(layer, prev_input, prev_delta);
        }
        else if(net.types[i] == MAXPOOL){
            maxpool_layer layer = *(maxpool_layer *)net.layers[i];

 src/network_gpu.c

@@ -87,7 +87,7 @@
        }
        if(net.types[i] == CONVOLUTIONAL){
            convolutional_layer layer = *(convolutional_layer *)net.layers[i];
            backward_convolutional_layer_gpu(layer, prev_delta);
            backward_convolutional_layer_gpu(layer, prev_input, prev_delta);
        }
        else if(net.types[i] == COST){
            cost_layer layer = *(cost_layer *)net.layers[i];

 src/opencl.c

@@ -88,7 +88,7 @@

    }
    int index = getpid()%num_devices;
    index = 1;
    index = 0;
    printf("%d rand, %d devices, %d index\n", getpid(), num_devices, index);
    info.device = devices[index];
    fprintf(stderr, "Found %d device(s)\n", num_devices);

 src/server.c

@@ -6,13 +6,22 @@
#include <netdb.h>

#include "server.h"
#include "connected_layer.h"

#define MESSAGESIZE 512
#define MESSAGESIZE 50012
#define NUMFLOATS ((MESSAGESIZE-12)/4)
#define SERVER_PORT 9876
#define CLIENT_PORT 9879
#define STR(x) #x
#define PARAMETER_SERVER localhost

typedef struct{
    int layer;
    int wob;
    int offset;
    float data[NUMFLOATS];
} message;

int socket_setup(int port)
{
    static int fd = 0;                         /* our socket */
@@ -42,27 +51,38 @@
    return fd;
}

void server_update()
void server_update(network net)
{
    int fd = socket_setup(SERVER_PORT);
    struct sockaddr_in remaddr;     /* remote address */
    socklen_t addrlen = sizeof(remaddr);            /* length of addresses */
    int recvlen;                    /* # bytes received */
    unsigned char buf[MESSAGESIZE];     /* receive buffer */
    message m;

    int count = 0;
    while(1){
    recvlen = recvfrom(fd, buf, MESSAGESIZE, 0, (struct sockaddr *)&remaddr, &addrlen);
    buf[recvlen] = 0;
    printf("received %d bytes\n", recvlen);
    printf("%s\n", buf);
        memcpy(&m, buf, recvlen);
        //printf("received %d bytes\n", recvlen);
        //printf("layer %d wob %d offset %d\n", m.layer, m.wob, m.offset);
        ++count;
        if(count % 100 == 0) printf("%d\n", count);
    }
    //printf("%s\n", buf);
}

void client_update()
void client_update(network net)
{
    int fd = socket_setup(CLIENT_PORT);
    struct hostent *hp;     /* host information */
    struct sockaddr_in servaddr;    /* server address */
    printf("%ld %ld\n", sizeof(message), MESSAGESIZE);
    char *my_message = "this is a test message";

    unsigned char buf[MESSAGESIZE];
    message m;

    /* fill in the server's address and data */
    memset((char*)&servaddr, 0, sizeof(servaddr));
    servaddr.sin_family = AF_INET;
@@ -78,7 +98,39 @@
    memcpy((void *)&servaddr.sin_addr, hp->h_addr_list[0], hp->h_length);

    /* send a message to the server */
    if (sendto(fd, my_message, strlen(my_message), 0, (struct sockaddr *)&servaddr, sizeof(servaddr)) < 0) {
    int i, j, k;
    for(i = 0; i < net.n; ++i){
        if(net.types[i] == CONNECTED){
            connected_layer *layer = (connected_layer *) net.layers[i];
            m.layer = i;
            m.wob = 0;
            for(j = 0; j < layer->outputs; j += NUMFLOATS){
                m.offset = j;

                int num = layer->outputs - j;
                if(NUMFLOATS < num) num = NUMFLOATS;

                memcpy(m.data, &layer->bias_updates[j], num*sizeof(float));
                memcpy(buf, &m, MESSAGESIZE);

                if (sendto(fd, buf, MESSAGESIZE, 0, (struct sockaddr *)&servaddr, sizeof(servaddr)) < 0) {
        perror("sendto failed");
    }
}
            m.wob = 1;
            for(j = 0; j < layer->outputs*layer->inputs; j += NUMFLOATS){
                m.offset = j;

                int num = layer->outputs*layer->inputs - j;
                if(NUMFLOATS < num) num = NUMFLOATS;

                memcpy(m.data, &layer->weight_updates[j], num*sizeof(float));
                memcpy(buf, &m, MESSAGESIZE);

                if (sendto(fd, buf, MESSAGESIZE, 0, (struct sockaddr *)&servaddr, sizeof(servaddr)) < 0) {
                    perror("sendto failed");
                }
            }
        }
    }
}

 src/server.h

@@ -1,3 +1,4 @@
#include "network.h"

void server_update();
void client_update();
void server_update(network net);
void client_update(network net);