mirror/darknet
1
0
Fork 0
mirror of https://github.com/pjreddie/darknet.git synced 2023-08-10 21:13:14 +03:00
Code Issues Projects Releases Wiki Activity

Detection good, split up col images

Browse Source
This commit is contained in:
Joseph Redmon 2014-12-03 23:20:29 -08:00
parent ff67f03476
commit 1edcf73a73
16 changed files with 282 additions and 221 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/axpy.cl
View File

@ -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]) ? 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)

74
src/cnn.c
View File

@ -36,6 +36,7 @@ void test_convolve()
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 @@ void test_convolutional_layer()
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_convolutional_layer()
*/
}
/*
void test_col2im()
{
float col[] = {1,2,1,2,
@ -116,13 +119,12 @@ void test_col2im()
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 @@ void test_col2im()
ksize, stride, pad, data_col) ;
for(i = 0; i < 18; ++i)printf("%f,", data_col[i]);
printf("\n");
*/
}
*/
#endif
@ -158,7 +160,7 @@ void test_convolve_matrix()
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 test_color()
void verify_convolutional_layer()
{
/*
srand(0);
int i;
int n = 1;
@ -225,6 +228,7 @@ void verify_convolutional_layer()
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 @@ void draw_detection(image im, float *box)
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 draw_detection(image im, float *box)
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 @@ void test_im2row()
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 @@ void test_gpu_net()
#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 @@ int main(int argc, char *argv[])
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();

23
src/col2im.c
View File

@ -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 @@ void col2im_cpu(float* data_col, int batch,
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 @@ void col2im_cpu(float* data_col, int batch,
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 @@ cl_kernel get_col2im_kernel()
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 @@ void col2im_ongpu(cl_mem data_col, int batch,
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 @@ void col2im_ongpu(cl_mem data_col, int batch,
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 @@ void col2im_gpu(float *data_col, int batch,
clReleaseMemObject(col_gpu);
clReleaseMemObject(im_gpu);
}
*/
#endif

15
src/col2im.cl
View File

@ -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 @@ __kernel void col2im(__global float *data_col, int batch,
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;
}

90
src/convolutional_layer.c
View File

@ -65,7 +65,7 @@ convolutional_layer *make_convolutional_layer(int batch, int h, int w, int c, in
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 @@ convolutional_layer *make_convolutional_layer(int batch, int h, int w, int c, in
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 @@ convolutional_layer *make_convolutional_layer(int batch, int h, int w, int c, in
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 @@ void resize_convolutional_layer(convolutional_layer *layer, int h, int w, int c)
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 @@ void forward_convolutional_layer(const convolutional_layer layer, float *in)
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 learn_bias_convolutional_layer(convolutional_layer layer)
}
}
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 @@ void backward_convolutional_layer(convolutional_layer layer, float *delta)
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 @@ void forward_convolutional_layer_gpu(convolutional_layer layer, cl_mem in)
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 @@ void backward_convolutional_layer_gpu(convolutional_layer layer, cl_mem delta_cl
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);
}
}
}

4
src/convolutional_layer.h
View File

@ -47,7 +47,7 @@ typedef struct {
#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 forward_convolutional_layer(const convolutional_layer layer, float *in);
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);

3
src/cost_layer.c
View File

@ -52,6 +52,7 @@ void forward_cost_layer(cost_layer layer, float *input, float *truth)
}
}
*(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 @@ void forward_cost_layer_gpu(cost_layer layer, cl_mem input, cl_mem truth)
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)

25
src/im2col.c
View File

@ -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 @@ void im2col_cpu(float* data_im, int batch,
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 @@ void im2col_cpu(float* data_im, int batch,
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 @@ cl_kernel get_im2col_nopad_kernel()
}
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 @@ void im2col_ongpu(cl_mem data_im, int batch,
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 @@ void im2col_ongpu(cl_mem data_im, int batch,
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 @@ void im2col_gpu(float *data_im, int batch,
clReleaseMemObject(col_gpu);
clReleaseMemObject(im_gpu);
}
*/
#endif

16
src/im2col.cl
View File

@ -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 @@ __kernel void im2col_pad(__global float *im, int batch,
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 @@ __kernel void im2col_pad(__global float *im, int batch,
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 @@ __kernel void im2col_nopad(__global float *im, int batch,
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 @@ __kernel void im2col_nopad(__global float *im, int batch,
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;

8
src/image.c
View File

@ -9,14 +9,14 @@ void draw_box(image a, int x1, int y1, int x2, int y2)
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;
}
}
}

10
src/mini_blas.h
View File

@ -15,18 +15,18 @@ void axpy_ongpu_offset(int N, float ALPHA, cl_mem X, int OFFX, int INCX, cl_mem
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 @@ void gemm_ongpu(int TA, int TB, int M, int N, int K, float ALPHA,
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);

2
src/network.c
View File

@ -213,7 +213,7 @@ void backward_network(network net, float *input)
}
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];

2
src/network_gpu.c
View File

@ -87,7 +87,7 @@ void backward_network_gpu(network net, cl_mem input)
}
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];

2
src/opencl.c
View File

@ -88,7 +88,7 @@ cl_info cl_init()
}
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);

66
src/server.c
View File

@ -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 @@ int socket_setup(int port)
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 @@ void client_update()
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");
}
}
}
}
}

5
src/server.h
View File

@ -1,3 +1,4 @@
#include "network.h"
void server_update();
void client_update();
void server_update(network net);
void client_update(network net);