From 1edcf73a73d2007afc61289245763f5cf0c29e10 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Thu, 04 Dec 2014 07:20:29 +0000
Subject: [PATCH] Detection good, split up col images
---
src/network.c | 2
src/network_gpu.c | 2
src/cost_layer.c | 3
src/server.h | 5
src/cnn.c | 74 ++++++-
src/im2col.c | 99 +++++-----
src/server.c | 70 ++++++-
src/convolutional_layer.h | 4
src/axpy.cl | 2
src/im2col.cl | 16 -
src/image.c | 8
src/convolutional_layer.c | 98 +++-------
src/col2im.c | 81 ++++----
src/mini_blas.h | 22 +-
src/col2im.cl | 15 -
src/opencl.c | 2
16 files changed, 282 insertions(+), 221 deletions(-)
diff --git a/src/axpy.cl b/src/axpy.cl
index 396389d..04eb534 100644
--- a/src/axpy.cl
+++ b/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)
diff --git a/src/cnn.c b/src/cnn.c
index a6627b1..46248ed 100644
--- a/src/cnn.c
+++ b/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();
diff --git a/src/col2im.c b/src/col2im.c
index f585226..c11f453 100644
--- a/src/col2im.c
+++ b/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,21 +24,18 @@
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;
- int c_im = c / ksize / ksize;
- for (h = 0; h < height_col; ++h) {
- 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;
- double val = data_col[col_index];
- col2im_add_pixel(data_im, height, width, channels,
- b, im_row, im_col, c_im, pad, val);
- }
+ for (c = 0; c < channels_col; ++c) {
+ int w_offset = c % ksize;
+ int h_offset = (c / ksize) % ksize;
+ int c_im = c / ksize / ksize;
+ for (h = 0; h < height_col; ++h) {
+ 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;
+ double val = data_col[col_index];
+ col2im_add_pixel(data_im, height, width, channels,
+ im_row, im_col, c_im, pad, val);
}
}
}
@@ -60,9 +57,9 @@
return im2col_kernel;
}
-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 col2im_ongpu(cl_mem data_col, int offset,
+ int channels, int height, int width,
+ int ksize, int stride, int pad, cl_mem data_im)
{
cl_setup();
cl_kernel kernel = get_col2im_kernel();
@@ -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,32 +77,34 @@
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)
-{
- int height_col = (height - ksize) / stride + 1;
- int width_col = (width - ksize) / stride + 1;
- int channels_col = channels * ksize * ksize;
+/*
+ void col2im_gpu(float *data_col, int batch,
+ int channels, int height, int width,
+ int ksize, int stride, int pad, float *data_im)
+ {
+ int height_col = (height - ksize) / stride + 1;
+ int width_col = (width - ksize) / stride + 1;
+ int channels_col = channels * ksize * ksize;
- size_t size = height_col*width_col*channels_col*batch;
- cl_mem col_gpu = cl_make_array(data_col, size);
- size = channels*height*width*batch;
- cl_mem im_gpu = cl_make_array(data_im, size);
+ size_t size = height_col*width_col*channels_col*batch;
+ cl_mem col_gpu = cl_make_array(data_col, size);
+ size = channels*height*width*batch;
+ cl_mem im_gpu = cl_make_array(data_im, size);
- col2im_ongpu(col_gpu, batch, channels, height, width,
- ksize, stride, pad, im_gpu);
+ col2im_ongpu(col_gpu, batch, channels, height, width,
+ ksize, stride, pad, im_gpu);
- cl_read_array(im_gpu, data_im, size);
- clReleaseMemObject(col_gpu);
- clReleaseMemObject(im_gpu);
-}
+ cl_read_array(im_gpu, data_im, size);
+ clReleaseMemObject(col_gpu);
+ clReleaseMemObject(im_gpu);
+ }
+ */
#endif
diff --git a/src/col2im.cl b/src/col2im.cl
index 2ccf89e..617e818 100644
--- a/src/col2im.cl
+++ b/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;
}
diff --git a/src/convolutional_layer.c b/src/convolutional_layer.c
index 77e6483..bae06d3 100644
--- a/src/convolutional_layer.c
+++ b/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;
- float *b = layer.col_image;
- float *c = layer.filter_updates;
+ 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;
+
+ 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);
+ if(delta){
+ a = layer.filters;
+ b = layer.delta + i*m*k;
+ c = layer.col_image;
- a = layer.filters;
- b = layer.delta;
- c = layer.col_image;
+ gemm(1,0,n,k,m,1,a,n,b,k,0,c,k);
- 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;
+ 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);
+ if(delta_cl){
- 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);
+ }
}
}
diff --git a/src/convolutional_layer.h b/src/convolutional_layer.h
index 970a9b1..1ceca16 100644
--- a/src/convolutional_layer.h
+++ b/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);
diff --git a/src/cost_layer.c b/src/cost_layer.c
index 6614b94..08d3bb5 100644
--- a/src/cost_layer.c
+++ b/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)
diff --git a/src/im2col.c b/src/im2col.c
index a56463b..18e6c0e 100644
--- a/src/im2col.c
+++ b/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,20 +26,17 @@
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;
- int c_im = c / ksize / ksize;
- for (h = 0; h < height_col; ++h) {
- 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;
- data_col[col_index] = im2col_get_pixel(data_im, height, width, channels,
- b, im_row, im_col, c_im, pad);
- }
+ for (c = 0; c < channels_col; ++c) {
+ int w_offset = c % ksize;
+ int h_offset = (c / ksize) % ksize;
+ int c_im = c / ksize / ksize;
+ for (h = 0; h < height_col; ++h) {
+ 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;
+ data_col[col_index] = im2col_get_pixel(data_im, height, width, channels,
+ im_row, im_col, c_im, pad);
}
}
}
@@ -74,9 +71,9 @@
}
-void im2col_ongpu(cl_mem data_im, int batch,
- int channels, int height, int width,
- int ksize, int stride, int pad, cl_mem data_col)
+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)
{
cl_setup();
@@ -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,45 +101,47 @@
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,
- int channels, int height, int width,
- int ksize, int stride, int pad, float *data_col)
-{
- cl_setup();
- cl_context context = cl.context;
- cl_command_queue queue = cl.queue;
+/*
+ void im2col_gpu(float *data_im,
+ int channels, int height, int width,
+ int ksize, int stride, int pad, float *data_col)
+ {
+ cl_setup();
+ cl_context context = cl.context;
+ cl_command_queue queue = cl.queue;
- size_t size = sizeof(float)*(channels*height*width*batch);
- cl_mem im_gpu = clCreateBuffer(context,
- CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
- size, data_im, &cl.error);
- check_error(cl);
+ size_t size = sizeof(float)*(channels*height*width*batch);
+ cl_mem im_gpu = clCreateBuffer(context,
+ CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
+ size, data_im, &cl.error);
+ check_error(cl);
- int height_col = (height - ksize) / stride + 1;
- int width_col = (width - ksize) / stride + 1;
- int channels_col = channels * ksize * ksize;
+ int height_col = (height - ksize) / stride + 1;
+ int width_col = (width - ksize) / stride + 1;
+ int channels_col = channels * ksize * ksize;
- size = sizeof(float)*(height_col*width_col*channels_col*batch);
- cl_mem col_gpu = clCreateBuffer(context,
- CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR,
- size, data_col, &cl.error);
- check_error(cl);
+ size = sizeof(float)*(height_col*width_col*channels_col*batch);
+ cl_mem col_gpu = clCreateBuffer(context,
+ CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR,
+ size, data_col, &cl.error);
+ check_error(cl);
- im2col_ongpu(im_gpu, batch, channels, height, width,
- ksize, stride, pad, col_gpu);
+ im2col_ongpu(im_gpu, batch, channels, height, width,
+ ksize, stride, pad, col_gpu);
- clEnqueueReadBuffer(queue, col_gpu, CL_TRUE, 0, size, data_col, 0, 0, 0);
- check_error(cl);
+ clEnqueueReadBuffer(queue, col_gpu, CL_TRUE, 0, size, data_col, 0, 0, 0);
+ check_error(cl);
- clReleaseMemObject(col_gpu);
- clReleaseMemObject(im_gpu);
-}
+ clReleaseMemObject(col_gpu);
+ clReleaseMemObject(im_gpu);
+ }
+ */
#endif
diff --git a/src/im2col.cl b/src/im2col.cl
index 0555274..31f7db6 100644
--- a/src/im2col.cl
+++ b/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;
diff --git a/src/image.c b/src/image.c
index 99837e1..b946be8 100644
--- a/src/image.c
+++ b/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;
}
}
}
diff --git a/src/mini_blas.h b/src/mini_blas.h
index 07b7cc6..ea73ed6 100644
--- a/src/mini_blas.h
+++ b/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,
- int channels, int height, int width,
- int ksize, int stride, int pad, float* data_col);
+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);
@@ -59,10 +59,10 @@
float BETA,
float *C, int ldc);
void gemm_cpu(int TA, int TB, int M, int N, int K, float ALPHA,
- float *A, int lda,
- float *B, int ldb,
- float BETA,
- float *C, int ldc);
+ float *A, int lda,
+ float *B, int ldb,
+ float BETA,
+ float *C, int ldc);
void axpy_cpu(int N, float ALPHA, float *X, int INCX, float *Y, int INCY);
void copy_cpu(int N, float *X, int INCX, float *Y, int INCY);
void scal_cpu(int N, float ALPHA, float *X, int INCX);
diff --git a/src/network.c b/src/network.c
index 339e6eb..3a6a184 100644
--- a/src/network.c
+++ b/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];
diff --git a/src/network_gpu.c b/src/network_gpu.c
index 938a014..fe53d0c 100644
--- a/src/network_gpu.c
+++ b/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];
diff --git a/src/opencl.c b/src/opencl.c
index 55fb56c..981067a 100644
--- a/src/opencl.c
+++ b/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);
diff --git a/src/server.c b/src/server.c
index 9837960..bcb59f5 100644
--- a/src/server.c
+++ b/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;
- recvlen = recvfrom(fd, buf, MESSAGESIZE, 0, (struct sockaddr *)&remaddr, &addrlen);
- buf[recvlen] = 0;
- printf("received %d bytes\n", recvlen);
- printf("%s\n", buf);
+ int count = 0;
+ while(1){
+ recvlen = recvfrom(fd, buf, MESSAGESIZE, 0, (struct sockaddr *)&remaddr, &addrlen);
+ 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) {
- perror("sendto failed");
+ 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");
+ }
+ }
+ }
}
}
diff --git a/src/server.h b/src/server.h
index bfffdaa..3d8a46b 100644
--- a/src/server.h
+++ b/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);
--
Gitblit v1.10.0