From 4ab366a805a7678642539465d68ef906b4599aeb Mon Sep 17 00:00:00 2001 From: Joseph Redmon Date: Mon, 22 Dec 2014 14:35:37 -0800 Subject: [PATCH] some fixes, some other experiments --- Makefile | 2 +- src/axpy.cl | 2 +- src/cnn.c | 38 +++++++++++++++++++++----------------- src/connected_layer.c | 36 ++++++++++++++++++++++++++++++++++-- src/connected_layer.h | 5 +++-- src/data.c | 42 ++++++++++++++++-------------------------- src/dropout_layer.c | 1 + src/image.c | 1 + src/network.c | 3 ++- src/network_gpu.c | 1 + 10 files changed, 81 insertions(+), 50 deletions(-) diff --git a/Makefile b/Makefile index a76c5327..32479995 100644 --- a/Makefile +++ b/Makefile @@ -27,7 +27,7 @@ LDFLAGS+= -lOpenCL endif endif CFLAGS= $(COMMON) $(OPTS) -CFLAGS= $(COMMON) -O0 -g +#CFLAGS= $(COMMON) -O0 -g LDFLAGS+=`pkg-config --libs opencv` -lm -pthread VPATH=./src/ EXEC=cnn diff --git a/src/axpy.cl b/src/axpy.cl index 04eb5343..1503e8f7 100644 --- a/src/axpy.cl +++ b/src/axpy.cl @@ -13,7 +13,7 @@ __kernel void scal(int N, float ALPHA, __global float *X, int INCX) __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] || i%mod == 0) ? x[i] : 0; + x[i] = (i%mod && !mask[(i/mod)*mod]) ? 0 : x[i]; } __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 fd83ee89..59948aae 100644 --- a/src/cnn.c +++ b/src/cnn.c @@ -31,21 +31,23 @@ void test_parser() save_network(net, "cfg/trained_imagenet_smaller.cfg"); } +#define AMNT 3 void draw_detection(image im, float *box, int side) { int j; int r, c; - float amount[5] = {0,0,0,0,0}; + float amount[AMNT] = {0}; for(r = 0; r < side*side; ++r){ - for(j = 0; j < 5; ++j){ - if(box[r*5] > amount[j]) { - amount[j] = box[r*5]; - break; + float val = box[r*5]; + for(j = 0; j < AMNT; ++j){ + if(val > amount[j]) { + float swap = val; + val = amount[j]; + amount[j] = swap; } } } - float smallest = amount[0]; - for(j = 1; j < 5; ++j) if(amount[j] < smallest) smallest = amount[j]; + float smallest = amount[AMNT-1]; for(r = 0; r < side; ++r){ for(c = 0; c < side; ++c){ @@ -57,9 +59,9 @@ void draw_detection(image im, float *box, int side) int x = c*d+box[j+2]*d; int h = box[j+3]*256; int w = box[j+4]*256; - printf("%f %f %f %f\n", box[j+1], box[j+2], box[j+3], box[j+4]); - printf("%d %d %d %d\n", x, y, w, h); - printf("%d %d %d %d\n", x-w/2, y-h/2, x+w/2, y+h/2); + //printf("%f %f %f %f\n", box[j+1], box[j+2], box[j+3], box[j+4]); + //printf("%d %d %d %d\n", x, y, w, h); + //printf("%d %d %d %d\n", x-w/2, y-h/2, x+w/2, y+h/2); draw_box(im, x-w/2, y-h/2, x+w/2, y+h/2); } } @@ -87,9 +89,11 @@ void train_detection_net() i += 1; time=clock(); data train = load_data_detection_jitter_random(imgs, paths, plist->size, 256, 256, 7, 7, 256); - /* - image im = float_to_image(224, 224, 3, train.X.vals[0]); - draw_detection(im, train.y.vals[0], 7); + //data train = load_data_detection_random(imgs, paths, plist->size, 224, 224, 7, 7, 256); + +/* + image im = float_to_image(224, 224, 3, train.X.vals[923]); + draw_detection(im, train.y.vals[923], 7); */ normalize_data_rows(train); @@ -151,10 +155,10 @@ void train_imagenet(char *cfgfile) //network net = parse_network_cfg("/home/pjreddie/imagenet_backup/alexnet_1270.cfg"); srand(time(0)); network net = parse_network_cfg(cfgfile); - set_learning_network(&net, net.learning_rate, .5, .0005); + set_learning_network(&net, net.learning_rate/10., .5, .0005); printf("Learning Rate: %g, Momentum: %g, Decay: %g\n", net.learning_rate, net.momentum, net.decay); int imgs = 1024; - int i = 23030; + int i = 44700; 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); @@ -385,8 +389,8 @@ void train_nist(char *cfgfile) data test = load_categorical_data_csv("data/mnist/mnist_test.csv",0,10); network net = parse_network_cfg(cfgfile); int count = 0; - int iters = 60000/net.batch + 1; - while(++count <= 10){ + int iters = 6000/net.batch + 1; + while(++count <= 100){ clock_t start = clock(), end; normalize_data_rows(train); normalize_data_rows(test); diff --git a/src/connected_layer.c b/src/connected_layer.c index 96236a34..938b8b86 100644 --- a/src/connected_layer.c +++ b/src/connected_layer.c @@ -24,15 +24,21 @@ connected_layer *make_connected_layer(int batch, int inputs, int outputs, ACTIVA layer->delta = calloc(batch*outputs, sizeof(float*)); layer->weight_updates = calloc(inputs*outputs, sizeof(float)); + layer->bias_updates = calloc(outputs, sizeof(float)); + + layer->weight_prev = calloc(inputs*outputs, sizeof(float)); + layer->bias_prev = calloc(outputs, sizeof(float)); + layer->weights = calloc(inputs*outputs, sizeof(float)); + layer->biases = calloc(outputs, sizeof(float)); + + float scale = 1./sqrt(inputs); //scale = .01; for(i = 0; i < inputs*outputs; ++i){ layer->weights[i] = scale*rand_normal(); } - layer->bias_updates = calloc(outputs, sizeof(float)); - layer->biases = calloc(outputs, sizeof(float)); for(i = 0; i < outputs; ++i){ layer->biases[i] = scale; } @@ -52,6 +58,32 @@ connected_layer *make_connected_layer(int batch, int inputs, int outputs, ACTIVA return layer; } +void secret_update_connected_layer(connected_layer *layer) +{ + int n = layer->outputs*layer->inputs; + float dot = dot_cpu(n, layer->weight_updates, 1, layer->weight_prev, 1); + float mag = sqrt(dot_cpu(n, layer->weight_updates, 1, layer->weight_updates, 1)) + * sqrt(dot_cpu(n, layer->weight_prev, 1, layer->weight_prev, 1)); + float cos = dot/mag; + if(cos > .3) layer->learning_rate *= 1.1; + else if (cos < -.3) layer-> learning_rate /= 1.1; + + scal_cpu(n, layer->momentum, layer->weight_prev, 1); + axpy_cpu(n, 1, layer->weight_updates, 1, layer->weight_prev, 1); + scal_cpu(n, 0, layer->weight_updates, 1); + + scal_cpu(layer->outputs, layer->momentum, layer->bias_prev, 1); + axpy_cpu(layer->outputs, 1, layer->bias_updates, 1, layer->bias_prev, 1); + scal_cpu(layer->outputs, 0, layer->bias_updates, 1); + + //printf("rate: %f\n", layer->learning_rate); + + axpy_cpu(layer->outputs, layer->learning_rate, layer->bias_prev, 1, layer->biases, 1); + + axpy_cpu(layer->inputs*layer->outputs, -layer->decay, layer->weights, 1, layer->weight_prev, 1); + axpy_cpu(layer->inputs*layer->outputs, layer->learning_rate, layer->weight_prev, 1, layer->weights, 1); +} + void update_connected_layer(connected_layer layer) { axpy_cpu(layer.outputs, layer.learning_rate, layer.bias_updates, 1, layer.biases, 1); diff --git a/src/connected_layer.h b/src/connected_layer.h index 1e5b4a70..0895728d 100644 --- a/src/connected_layer.h +++ b/src/connected_layer.h @@ -18,8 +18,8 @@ typedef struct{ float *weight_updates; float *bias_updates; - float *weight_adapt; - float *bias_adapt; + float *weight_prev; + float *bias_prev; float *output; float *delta; @@ -38,6 +38,7 @@ typedef struct{ } connected_layer; +void secret_update_connected_layer(connected_layer *layer); connected_layer *make_connected_layer(int batch, int inputs, int outputs, ACTIVATION activation, float learning_rate, float momentum, float decay); void forward_connected_layer(connected_layer layer, float *input); diff --git a/src/data.c b/src/data.c index 86e59efc..3f74f6bd 100644 --- a/src/data.c +++ b/src/data.c @@ -81,6 +81,18 @@ matrix load_image_paths(char **paths, int n, int h, int w) return X; } +char **get_random_paths(char **paths, int n, int m) +{ + char **random_paths = calloc(n, sizeof(char*)); + int i; + for(i = 0; i < n; ++i){ + int index = rand()%m; + random_paths[i] = paths[index]; + if(i == 0) printf("%s\n", paths[index]); + } + return random_paths; +} + matrix load_labels_paths(char **paths, int n, char **labels, int k) { matrix y = make_matrix(n, k); @@ -138,13 +150,8 @@ void free_data(data d) data load_data_detection_jitter_random(int n, char **paths, int m, int h, int w, int nh, int nw, float scale) { - char **random_paths = calloc(n, sizeof(char*)); + char **random_paths = get_random_paths(paths, n, m); int i; - for(i = 0; i < n; ++i){ - int index = rand()%m; - random_paths[i] = paths[index]; - if(i == 0) printf("%s\n", paths[index]); - } data d; d.shallow = 0; d.X = load_image_paths(random_paths, n, h, w); @@ -154,10 +161,11 @@ data load_data_detection_jitter_random(int n, char **paths, int m, int h, int w, int dx = rand()%32; int dy = rand()%32; fill_truth_detection(random_paths[i], d.y.vals[i], 224, 224, nh, nw, scale, dx, dy); - image a = float_to_image(h, w, 3, d.X.vals[i]); jitter_image(a,224,224,dy,dx); } + d.X.cols = 224*224*3; + // print_matrix(d.y); free(random_paths); return d; } @@ -165,13 +173,7 @@ data load_data_detection_jitter_random(int n, char **paths, int m, int h, int w, data load_data_detection_random(int n, char **paths, int m, int h, int w, int nh, int nw, float scale) { - char **random_paths = calloc(n, sizeof(char*)); - int i; - for(i = 0; i < n; ++i){ - int index = rand()%m; - random_paths[i] = paths[index]; - if(i == 0) printf("%s\n", paths[index]); - } + char **random_paths = get_random_paths(paths, n, m); data d; d.shallow = 0; d.X = load_image_paths(random_paths, n, h, w); @@ -180,18 +182,6 @@ data load_data_detection_random(int n, char **paths, int m, int h, int w, int nh return d; } -char **get_random_paths(char **paths, int n, int m) -{ - char **random_paths = calloc(n, sizeof(char*)); - int i; - for(i = 0; i < n; ++i){ - int index = rand()%m; - random_paths[i] = paths[index]; - if(i == 0) printf("%s\n", paths[index]); - } - return random_paths; -} - data load_data(char **paths, int n, int m, char **labels, int k, int h, int w) { if(m) paths = get_random_paths(paths, n, m); diff --git a/src/dropout_layer.c b/src/dropout_layer.c index 8104b56f..edcb4268 100644 --- a/src/dropout_layer.c +++ b/src/dropout_layer.c @@ -80,6 +80,7 @@ void forward_dropout_layer_gpu(dropout_layer layer, cl_mem input) void backward_dropout_layer_gpu(dropout_layer layer, cl_mem delta) { + if(!delta) return; int size = layer.inputs*layer.batch; cl_kernel kernel = get_dropout_kernel(); diff --git a/src/image.c b/src/image.c index a2664a9b..ddb5bf52 100644 --- a/src/image.c +++ b/src/image.c @@ -39,6 +39,7 @@ void jitter_image(image a, int h, int w, int dh, int dw) for(j = 0; j < w; ++j){ int src = j + dw + (i+dh)*a.w + k*a.w*a.h; int dst = j + i*w + k*w*h; + //printf("%d %d\n", src, dst); a.data[dst] = a.data[src]; } } diff --git a/src/network.c b/src/network.c index 0bf5357c..42253dc9 100644 --- a/src/network.c +++ b/src/network.c @@ -103,7 +103,8 @@ void update_network(network net) } else if(net.types[i] == CONNECTED){ connected_layer layer = *(connected_layer *)net.layers[i]; - update_connected_layer(layer); + secret_update_connected_layer((connected_layer *)net.layers[i]); + //update_connected_layer(layer); } } } diff --git a/src/network_gpu.c b/src/network_gpu.c index 6ff95c8c..4d2c8d34 100644 --- a/src/network_gpu.c +++ b/src/network_gpu.c @@ -195,6 +195,7 @@ float *get_network_output_layer_gpu(network net, int i) } else if(net.types[i] == CONNECTED){ connected_layer layer = *(connected_layer *)net.layers[i]; + cl_read_array(layer.output_cl, layer.output, layer.outputs*layer.batch); return layer.output; } else if(net.types[i] == MAXPOOL){