From 913d355ec1cf34aad71fdd75202fc3b0309e63a0 Mon Sep 17 00:00:00 2001 From: Joseph Redmon Date: Thu, 28 Jan 2016 12:30:38 -0800 Subject: [PATCH] lots of stuff --- Makefile | 2 +- src/blas.c | 12 +- src/blas.h | 11 ++ src/blas_kernels.cu | 16 ++ src/connected_layer.c | 151 +++++++++++++++++-- src/connected_layer.h | 2 +- src/convolutional_kernels.cu | 42 +++++- src/convolutional_layer.c | 71 ++++++++- src/convolutional_layer.h | 3 +- src/cost_layer.c | 21 ++- src/crop_layer.h | 1 - src/crop_layer_kernels.cu | 2 +- src/darknet.c | 8 +- src/deconvolutional_layer.h | 1 - src/detection_layer.c | 2 +- src/dropout_layer.c | 1 - src/dropout_layer.h | 1 - src/dropout_layer_kernels.cu | 1 - src/image.c | 2 + src/layer.h | 26 +++- src/maxpool_layer.h | 1 - src/network.c | 9 ++ src/network.h | 3 +- src/network_kernels.cu | 8 +- src/nightmare.c | 27 +++- src/params.h | 1 - src/parser.c | 98 ++++++++++--- src/rnn.c | 147 +++++++++++++++++++ src/rnn_layer.c | 275 +++++++++++++++++++++++++++++++++++ src/rnn_layer.h | 24 +++ src/softmax_layer.c | 10 +- src/softmax_layer.h | 3 +- src/softmax_layer_kernels.cu | 10 +- src/utils.c | 3 +- src/yolo.c | 4 +- 35 files changed, 913 insertions(+), 86 deletions(-) delete mode 100644 src/params.h create mode 100644 src/rnn.c create mode 100644 src/rnn_layer.c create mode 100644 src/rnn_layer.h diff --git a/Makefile b/Makefile index 04e30f35..c9b6ecac 100644 --- a/Makefile +++ b/Makefile @@ -34,7 +34,7 @@ CFLAGS+= -DGPU LDFLAGS+= -L/usr/local/cuda/lib64 -lcuda -lcudart -lcublas -lcurand endif -OBJ=gemm.o utils.o cuda.o deconvolutional_layer.o convolutional_layer.o list.o image.o activations.o im2col.o col2im.o blas.o crop_layer.o dropout_layer.o maxpool_layer.o softmax_layer.o data.o matrix.o network.o connected_layer.o cost_layer.o parser.o option_list.o darknet.o detection_layer.o imagenet.o captcha.o route_layer.o writing.o box.o nightmare.o normalization_layer.o avgpool_layer.o coco.o dice.o yolo.o layer.o compare.o classifier.o local_layer.o swag.o shortcut_layer.o activation_layer.o +OBJ=gemm.o utils.o cuda.o deconvolutional_layer.o convolutional_layer.o list.o image.o activations.o im2col.o col2im.o blas.o crop_layer.o dropout_layer.o maxpool_layer.o softmax_layer.o data.o matrix.o network.o connected_layer.o cost_layer.o parser.o option_list.o darknet.o detection_layer.o imagenet.o captcha.o route_layer.o writing.o box.o nightmare.o normalization_layer.o avgpool_layer.o coco.o dice.o yolo.o layer.o compare.o classifier.o local_layer.o swag.o shortcut_layer.o activation_layer.o rnn_layer.o rnn.o ifeq ($(GPU), 1) OBJ+=convolutional_kernels.o deconvolutional_kernels.o activation_kernels.o im2col_kernels.o col2im_kernels.o blas_kernels.o crop_layer_kernels.o dropout_layer_kernels.o maxpool_layer_kernels.o softmax_layer_kernels.o network_kernels.o avgpool_layer_kernels.o yolo_kernels.o coco_kernels.o endif diff --git a/src/blas.c b/src/blas.c index 8769df35..d7948bb1 100644 --- a/src/blas.c +++ b/src/blas.c @@ -67,7 +67,7 @@ void normalize_cpu(float *x, float *mean, float *variance, int batch, int filter for(f = 0; f < filters; ++f){ for(i = 0; i < spatial; ++i){ int index = b*filters*spatial + f*spatial + i; - x[index] = (x[index] - mean[f])/(sqrt(variance[f])); + x[index] = (x[index] - mean[f])/(sqrt(variance[f]) + .00001f); } } } @@ -115,6 +115,16 @@ void copy_cpu(int N, float *X, int INCX, float *Y, int INCY) for(i = 0; i < N; ++i) Y[i*INCY] = X[i*INCX]; } +void smooth_l1_cpu(int n, float *pred, float *truth, float *delta) +{ + int i; + for(i = 0; i < n; ++i){ + float diff = truth[i] - pred[i]; + if(fabs(diff) > 1) delta[i] = diff; + else delta[i] = (diff > 0) ? 1 : -1; + } +} + float dot_cpu(int N, float *X, int INCX, float *Y, int INCY) { int i; diff --git a/src/blas.h b/src/blas.h index aecdc593..f5189e5e 100644 --- a/src/blas.h +++ b/src/blas.h @@ -17,11 +17,18 @@ void fill_cpu(int N, float ALPHA, float * X, int INCX); float dot_cpu(int N, float *X, int INCX, float *Y, int INCY); void test_gpu_blas(); void shortcut_cpu(int batch, int w1, int h1, int c1, float *add, int w2, int h2, int c2, float *out); +void smooth_l1_cpu(int n, float *pred, float *truth, float *delta); void mean_cpu(float *x, int batch, int filters, int spatial, float *mean); void variance_cpu(float *x, float *mean, int batch, int filters, int spatial, float *variance); void normalize_cpu(float *x, float *mean, float *variance, int batch, int filters, int spatial); +void scale_bias(float *output, float *scales, int batch, int n, int size); +void backward_scale_cpu(float *x_norm, float *delta, int batch, int n, int size, float *scale_updates); +void mean_delta_cpu(float *delta, float *variance, int batch, int filters, int spatial, float *mean_delta); +void variance_delta_cpu(float *x, float *delta, float *mean, float *variance, int batch, int filters, int spatial, float *variance_delta); +void normalize_delta_cpu(float *x, float *mean, float *variance, float *mean_delta, float *variance_delta, int batch, int filters, int spatial, float *delta); + #ifdef GPU void axpy_ongpu(int N, float ALPHA, float * X, int INCX, float * Y, int INCY); void axpy_ongpu_offset(int N, float ALPHA, float * X, int OFFX, int INCX, float * Y, int OFFY, int INCY); @@ -46,5 +53,9 @@ void fast_variance_delta_gpu(float *x, float *delta, float *mean, float *varianc void fast_variance_gpu(float *x, float *mean, int batch, int filters, int spatial, float *variance); void fast_mean_gpu(float *x, int batch, int filters, int spatial, float *mean); void shortcut_gpu(int batch, int w1, int h1, int c1, float *add, int w2, int h2, int c2, float *out); +void smooth_l1_gpu(int n, float *pred, float *truth, float *delta); +void scale_bias_gpu(float *output, float *biases, int batch, int n, int size); +void backward_scale_gpu(float *x_norm, float *delta, int batch, int n, int size, float *scale_updates); +void scale_bias_gpu(float *output, float *biases, int batch, int n, int size); #endif #endif diff --git a/src/blas_kernels.cu b/src/blas_kernels.cu index 49406db2..61db29f2 100644 --- a/src/blas_kernels.cu +++ b/src/blas_kernels.cu @@ -409,3 +409,19 @@ extern "C" void shortcut_gpu(int batch, int w1, int h1, int c1, float *add, int shortcut_kernel<<>>(size, minw, minh, minc, stride, sample, batch, w1, h1, c1, add, w2, h2, c2, out); check_error(cudaPeekAtLastError()); } + +__global__ void smooth_l1_kernel(int n, float *pred, float *truth, float *delta) +{ + int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; + if(i < n){ + float diff = truth[i] - pred[i]; + if(abs(diff) > 1) delta[i] = diff; + else delta[i] = (diff > 0) ? 1 : -1; + } +} + +extern "C" void smooth_l1_gpu(int n, float *pred, float *truth, float *delta) +{ + smooth_l1_kernel<<>>(n, pred, truth, delta); + check_error(cudaPeekAtLastError()); +} diff --git a/src/connected_layer.c b/src/connected_layer.c index c0a9d8b3..df78e67b 100644 --- a/src/connected_layer.c +++ b/src/connected_layer.c @@ -9,7 +9,7 @@ #include #include -connected_layer make_connected_layer(int batch, int inputs, int outputs, ACTIVATION activation) +connected_layer make_connected_layer(int batch, int inputs, int outputs, ACTIVATION activation, int batch_normalize) { int i; connected_layer l = {0}; @@ -18,9 +18,10 @@ connected_layer make_connected_layer(int batch, int inputs, int outputs, ACTIVAT l.inputs = inputs; l.outputs = outputs; l.batch=batch; + l.batch_normalize = batch_normalize; - l.output = calloc(batch*outputs, sizeof(float*)); - l.delta = calloc(batch*outputs, sizeof(float*)); + l.output = calloc(batch*outputs, sizeof(float)); + l.delta = calloc(batch*outputs, sizeof(float)); l.weight_updates = calloc(inputs*outputs, sizeof(float)); l.bias_updates = calloc(outputs, sizeof(float)); @@ -39,6 +40,25 @@ connected_layer make_connected_layer(int batch, int inputs, int outputs, ACTIVAT l.biases[i] = scale; } + if(batch_normalize){ + l.scales = calloc(outputs, sizeof(float)); + l.scale_updates = calloc(outputs, sizeof(float)); + for(i = 0; i < outputs; ++i){ + l.scales[i] = 1; + } + + l.mean = calloc(outputs, sizeof(float)); + l.mean_delta = calloc(outputs, sizeof(float)); + l.variance = calloc(outputs, sizeof(float)); + l.variance_delta = calloc(outputs, sizeof(float)); + + l.rolling_mean = calloc(outputs, sizeof(float)); + l.rolling_variance = calloc(outputs, sizeof(float)); + + l.x = calloc(batch*outputs, sizeof(float)); + l.x_norm = calloc(batch*outputs, sizeof(float)); + } + #ifdef GPU l.weights_gpu = cuda_make_array(l.weights, outputs*inputs); l.biases_gpu = cuda_make_array(l.biases, outputs); @@ -48,6 +68,22 @@ connected_layer make_connected_layer(int batch, int inputs, int outputs, ACTIVAT l.output_gpu = cuda_make_array(l.output, outputs*batch); l.delta_gpu = cuda_make_array(l.delta, outputs*batch); + if(batch_normalize){ + l.scales_gpu = cuda_make_array(l.scales, outputs); + l.scale_updates_gpu = cuda_make_array(l.scale_updates, outputs); + + l.mean_gpu = cuda_make_array(l.mean, outputs); + l.variance_gpu = cuda_make_array(l.variance, outputs); + + l.rolling_mean_gpu = cuda_make_array(l.mean, outputs); + l.rolling_variance_gpu = cuda_make_array(l.variance, outputs); + + l.mean_delta_gpu = cuda_make_array(l.mean, outputs); + l.variance_delta_gpu = cuda_make_array(l.variance, outputs); + + l.x_gpu = cuda_make_array(l.output, l.batch*outputs); + l.x_norm_gpu = cuda_make_array(l.output, l.batch*outputs); + } #endif l.activation = activation; fprintf(stderr, "Connected Layer: %d inputs, %d outputs\n", inputs, outputs); @@ -59,6 +95,11 @@ void update_connected_layer(connected_layer l, int batch, float learning_rate, f axpy_cpu(l.outputs, learning_rate/batch, l.bias_updates, 1, l.biases, 1); scal_cpu(l.outputs, momentum, l.bias_updates, 1); + if(l.batch_normalize){ + axpy_cpu(l.outputs, learning_rate/batch, l.scale_updates, 1, l.scales, 1); + scal_cpu(l.outputs, momentum, l.scale_updates, 1); + } + axpy_cpu(l.inputs*l.outputs, -decay*batch, l.weights, 1, l.weight_updates, 1); axpy_cpu(l.inputs*l.outputs, learning_rate/batch, l.weight_updates, 1, l.weights, 1); scal_cpu(l.inputs*l.outputs, momentum, l.weight_updates, 1); @@ -67,9 +108,7 @@ void update_connected_layer(connected_layer l, int batch, float learning_rate, f void forward_connected_layer(connected_layer l, network_state state) { int i; - for(i = 0; i < l.batch; ++i){ - copy_cpu(l.outputs, l.biases, 1, l.output + i*l.outputs, 1); - } + fill_cpu(l.outputs*l.batch, 0, l.output, 1); int m = l.batch; int k = l.inputs; int n = l.outputs; @@ -77,6 +116,27 @@ void forward_connected_layer(connected_layer l, network_state state) float *b = l.weights; float *c = l.output; gemm(0,1,m,n,k,1,a,k,b,k,1,c,n); + if(l.batch_normalize){ + if(state.train){ + mean_cpu(l.output, l.batch, l.outputs, 1, l.mean); + variance_cpu(l.output, l.mean, l.batch, l.outputs, 1, l.variance); + + scal_cpu(l.outputs, .95, l.rolling_mean, 1); + axpy_cpu(l.outputs, .05, l.mean, 1, l.rolling_mean, 1); + scal_cpu(l.outputs, .95, l.rolling_variance, 1); + axpy_cpu(l.outputs, .05, l.variance, 1, l.rolling_variance, 1); + + copy_cpu(l.outputs*l.batch, l.output, 1, l.x, 1); + normalize_cpu(l.output, l.mean, l.variance, l.batch, l.outputs, 1); + copy_cpu(l.outputs*l.batch, l.output, 1, l.x_norm, 1); + } else { + normalize_cpu(l.output, l.rolling_mean, l.rolling_variance, l.batch, l.outputs, 1); + } + scale_bias(l.output, l.scales, l.batch, l.outputs, 1); + } + for(i = 0; i < l.batch; ++i){ + axpy_cpu(l.outputs, 1, l.biases, 1, l.output + i*l.outputs, 1); + } activate_array(l.output, l.outputs*l.batch, l.activation); } @@ -87,6 +147,16 @@ void backward_connected_layer(connected_layer l, network_state state) for(i = 0; i < l.batch; ++i){ axpy_cpu(l.outputs, 1, l.delta + i*l.outputs, 1, l.bias_updates, 1); } + if(l.batch_normalize){ + backward_scale_cpu(l.x_norm, l.delta, l.batch, l.outputs, 1, l.scale_updates); + + scale_bias(l.delta, l.scales, l.batch, l.outputs, 1); + + mean_delta_cpu(l.delta, l.variance, l.batch, l.outputs, 1, l.mean_delta); + variance_delta_cpu(l.x, l.delta, l.mean, l.variance, l.batch, l.outputs, 1, l.variance_delta); + normalize_delta_cpu(l.x, l.mean, l.variance, l.mean_delta, l.variance_delta, l.batch, l.outputs, 1, l.delta); + } + int m = l.outputs; int k = l.batch; int n = l.inputs; @@ -114,6 +184,11 @@ void pull_connected_layer(connected_layer l) cuda_pull_array(l.biases_gpu, l.biases, l.outputs); cuda_pull_array(l.weight_updates_gpu, l.weight_updates, l.inputs*l.outputs); cuda_pull_array(l.bias_updates_gpu, l.bias_updates, l.outputs); + if (l.batch_normalize){ + cuda_pull_array(l.scales_gpu, l.scales, l.outputs); + cuda_pull_array(l.rolling_mean_gpu, l.rolling_mean, l.outputs); + cuda_pull_array(l.rolling_variance_gpu, l.rolling_variance, l.outputs); + } } void push_connected_layer(connected_layer l) @@ -122,6 +197,11 @@ void push_connected_layer(connected_layer l) cuda_push_array(l.biases_gpu, l.biases, l.outputs); cuda_push_array(l.weight_updates_gpu, l.weight_updates, l.inputs*l.outputs); cuda_push_array(l.bias_updates_gpu, l.bias_updates, l.outputs); + if (l.batch_normalize){ + cuda_push_array(l.scales_gpu, l.scales, l.outputs); + cuda_push_array(l.rolling_mean_gpu, l.rolling_mean, l.outputs); + cuda_push_array(l.rolling_variance_gpu, l.rolling_variance, l.outputs); + } } void update_connected_layer_gpu(connected_layer l, int batch, float learning_rate, float momentum, float decay) @@ -129,6 +209,11 @@ void update_connected_layer_gpu(connected_layer l, int batch, float learning_rat axpy_ongpu(l.outputs, learning_rate/batch, l.bias_updates_gpu, 1, l.biases_gpu, 1); scal_ongpu(l.outputs, momentum, l.bias_updates_gpu, 1); + if(l.batch_normalize){ + axpy_ongpu(l.outputs, learning_rate/batch, l.scale_updates_gpu, 1, l.scales_gpu, 1); + scal_ongpu(l.outputs, momentum, l.scale_updates_gpu, 1); + } + axpy_ongpu(l.inputs*l.outputs, -decay*batch, l.weights_gpu, 1, l.weight_updates_gpu, 1); axpy_ongpu(l.inputs*l.outputs, learning_rate/batch, l.weight_updates_gpu, 1, l.weights_gpu, 1); scal_ongpu(l.inputs*l.outputs, momentum, l.weight_updates_gpu, 1); @@ -137,9 +222,12 @@ void update_connected_layer_gpu(connected_layer l, int batch, float learning_rat void forward_connected_layer_gpu(connected_layer l, network_state state) { int i; - for(i = 0; i < l.batch; ++i){ - copy_ongpu_offset(l.outputs, l.biases_gpu, 0, 1, l.output_gpu, i*l.outputs, 1); - } + fill_ongpu(l.outputs*l.batch, 0, l.output_gpu, 1); + /* + for(i = 0; i < l.batch; ++i){ + copy_ongpu_offset(l.outputs, l.biases_gpu, 0, 1, l.output_gpu, i*l.outputs, 1); + } + */ int m = l.batch; int k = l.inputs; int n = l.outputs; @@ -147,13 +235,35 @@ void forward_connected_layer_gpu(connected_layer l, network_state state) float * b = l.weights_gpu; float * c = l.output_gpu; gemm_ongpu(0,1,m,n,k,1,a,k,b,k,1,c,n); + if(l.batch_normalize){ + if(state.train){ + fast_mean_gpu(l.output_gpu, l.batch, l.outputs, 1, l.mean_gpu); + fast_variance_gpu(l.output_gpu, l.mean_gpu, l.batch, l.outputs, 1, l.variance_gpu); + + scal_ongpu(l.outputs, .95, l.rolling_mean_gpu, 1); + axpy_ongpu(l.outputs, .05, l.mean_gpu, 1, l.rolling_mean_gpu, 1); + scal_ongpu(l.outputs, .95, l.rolling_variance_gpu, 1); + axpy_ongpu(l.outputs, .05, l.variance_gpu, 1, l.rolling_variance_gpu, 1); + + copy_ongpu(l.outputs*l.batch, l.output_gpu, 1, l.x_gpu, 1); + normalize_gpu(l.output_gpu, l.mean_gpu, l.variance_gpu, l.batch, l.outputs, 1); + copy_ongpu(l.outputs*l.batch, l.output_gpu, 1, l.x_norm_gpu, 1); + } else { + normalize_gpu(l.output_gpu, l.rolling_mean_gpu, l.rolling_variance_gpu, l.batch, l.outputs, 1); + } + + scale_bias_gpu(l.output_gpu, l.scales_gpu, l.batch, l.outputs, 1); + } + for(i = 0; i < l.batch; ++i){ + axpy_ongpu(l.outputs, 1, l.biases_gpu, 1, l.output_gpu + i*l.outputs, 1); + } activate_array_ongpu(l.output_gpu, l.outputs*l.batch, l.activation); -/* - cuda_pull_array(l.output_gpu, l.output, l.outputs*l.batch); - float avg = mean_array(l.output, l.outputs*l.batch); - printf("%f\n", avg); - */ + /* + cuda_pull_array(l.output_gpu, l.output, l.outputs*l.batch); + float avg = mean_array(l.output, l.outputs*l.batch); + printf("%f\n", avg); + */ } void backward_connected_layer_gpu(connected_layer l, network_state state) @@ -161,8 +271,19 @@ void backward_connected_layer_gpu(connected_layer l, network_state state) int i; gradient_array_ongpu(l.output_gpu, l.outputs*l.batch, l.activation, l.delta_gpu); for(i = 0; i < l.batch; ++i){ - axpy_ongpu_offset(l.outputs, 1, l.delta_gpu, i*l.outputs, 1, l.bias_updates_gpu, 0, 1); + axpy_ongpu(l.outputs, 1, l.delta_gpu + i*l.outputs, 1, l.bias_updates_gpu, 1); } + + if(l.batch_normalize){ + backward_scale_gpu(l.x_norm_gpu, l.delta_gpu, l.batch, l.outputs, 1, l.scale_updates_gpu); + + scale_bias_gpu(l.delta_gpu, l.scales_gpu, l.batch, l.outputs, 1); + + fast_mean_delta_gpu(l.delta_gpu, l.variance_gpu, l.batch, l.outputs, 1, l.mean_delta_gpu); + fast_variance_delta_gpu(l.x_gpu, l.delta_gpu, l.mean_gpu, l.variance_gpu, l.batch, l.outputs, 1, l.variance_delta_gpu); + normalize_delta_gpu(l.x_gpu, l.mean_gpu, l.variance_gpu, l.mean_delta_gpu, l.variance_delta_gpu, l.batch, l.outputs, 1, l.delta_gpu); + } + int m = l.outputs; int k = l.batch; int n = l.inputs; diff --git a/src/connected_layer.h b/src/connected_layer.h index 2bf53b26..56bd1c38 100644 --- a/src/connected_layer.h +++ b/src/connected_layer.h @@ -7,7 +7,7 @@ typedef layer connected_layer; -connected_layer make_connected_layer(int batch, int inputs, int outputs, ACTIVATION activation); +connected_layer make_connected_layer(int batch, int inputs, int outputs, ACTIVATION activation, int batch_normalize); void forward_connected_layer(connected_layer layer, network_state state); void backward_connected_layer(connected_layer layer, network_state state); diff --git a/src/convolutional_kernels.cu b/src/convolutional_kernels.cu index a64a499e..4fdc1a1e 100644 --- a/src/convolutional_kernels.cu +++ b/src/convolutional_kernels.cu @@ -12,6 +12,21 @@ extern "C" { #include "cuda.h" } +__global__ void binarize_filters_kernel(float *filters, int n, int size, float *binary) +{ + int f = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; + if (f >= n) return; + int i = 0; + float mean = 0; + for(i = 0; i < size; ++i){ + mean += abs(filters[f*size + i]); + } + mean = mean / size; + for(i = 0; i < size; ++i){ + binary[f*size + i] = (filters[f*size + i] > 0) ? mean : -mean; + } +} + __global__ void scale_bias_kernel(float *output, float *biases, int n, int size) { int offset = blockIdx.x * blockDim.x + threadIdx.x; @@ -50,6 +65,12 @@ __global__ void backward_scale_kernel(float *x_norm, float *delta, int batch, in } } +void binarize_filters_gpu(float *filters, int n, int size, float *mean) +{ + binarize_filters_kernel<<>>(filters, n, size, mean); + check_error(cudaPeekAtLastError()); +} + void backward_scale_gpu(float *x_norm, float *delta, int batch, int n, int size, float *scale_updates) { backward_scale_kernel<<>>(x_norm, delta, batch, n, size, scale_updates); @@ -100,6 +121,13 @@ void backward_bias_gpu(float *bias_updates, float *delta, int batch, int n, int check_error(cudaPeekAtLastError()); } +void swap_binary(convolutional_layer l) +{ + float *swap = l.filters_gpu; + l.filters_gpu = l.binary_filters_gpu; + l.binary_filters_gpu = swap; +} + void forward_convolutional_layer_gpu(convolutional_layer l, network_state state) { int i; @@ -109,6 +137,11 @@ void forward_convolutional_layer_gpu(convolutional_layer l, network_state state) convolutional_out_width(l); fill_ongpu(l.outputs*l.batch, 0, l.output_gpu, 1); + if(l.binary){ + binarize_filters_gpu(l.filters_gpu, l.n, l.c*l.size*l.size, l.binary_filters_gpu); + swap_binary(l); + } + for(i = 0; i < l.batch; ++i){ im2col_ongpu(state.input + i*l.c*l.h*l.w, l.c, l.h, l.w, l.size, l.stride, l.pad, l.col_image_gpu); float * a = l.filters_gpu; @@ -122,12 +155,6 @@ void forward_convolutional_layer_gpu(convolutional_layer l, network_state state) fast_mean_gpu(l.output_gpu, l.batch, l.n, l.out_h*l.out_w, l.mean_gpu); fast_variance_gpu(l.output_gpu, l.mean_gpu, l.batch, l.n, l.out_h*l.out_w, l.variance_gpu); - /* - cuda_pull_array(l.variance_gpu, l.mean, 1); - printf("%f\n", l.mean[0]); - */ - - scal_ongpu(l.n, .95, l.rolling_mean_gpu, 1); axpy_ongpu(l.n, .05, l.mean_gpu, 1, l.rolling_mean_gpu, 1); scal_ongpu(l.n, .95, l.rolling_variance_gpu, 1); @@ -145,6 +172,7 @@ void forward_convolutional_layer_gpu(convolutional_layer l, network_state state) add_bias_gpu(l.output_gpu, l.biases_gpu, l.batch, l.n, n); activate_array_ongpu(l.output_gpu, m*n*l.batch, l.activation); + if(l.binary) swap_binary(l); } void backward_convolutional_layer_gpu(convolutional_layer l, network_state state) @@ -178,6 +206,7 @@ void backward_convolutional_layer_gpu(convolutional_layer l, network_state state gemm_ongpu(0,1,m,n,k,1,a + i*m*k,k,b,k,1,c,n); if(state.delta){ + if(l.binary) swap_binary(l); float * a = l.filters_gpu; float * b = l.delta_gpu; float * c = l.col_image_gpu; @@ -185,6 +214,7 @@ void backward_convolutional_layer_gpu(convolutional_layer l, network_state state gemm_ongpu(1,0,n,k,m,1,a,n,b + i*k*m,k,0,c,k); col2im_ongpu(l.col_image_gpu, l.c, l.h, l.w, l.size, l.stride, l.pad, state.delta + i*l.c*l.h*l.w); + if(l.binary) swap_binary(l); } } } diff --git a/src/convolutional_layer.c b/src/convolutional_layer.c index 871a84e1..159951dc 100644 --- a/src/convolutional_layer.c +++ b/src/convolutional_layer.c @@ -41,7 +41,65 @@ image get_convolutional_delta(convolutional_layer l) return float_to_image(w,h,c,l.delta); } -convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, int pad, ACTIVATION activation, int batch_normalize) +void backward_scale_cpu(float *x_norm, float *delta, int batch, int n, int size, float *scale_updates) +{ + int i,b,f; + for(f = 0; f < n; ++f){ + float sum = 0; + for(b = 0; b < batch; ++b){ + for(i = 0; i < size; ++i){ + int index = i + size*(f + n*b); + sum += delta[index] * x_norm[index]; + } + } + scale_updates[f] += sum; + } +} + +void mean_delta_cpu(float *delta, float *variance, int batch, int filters, int spatial, float *mean_delta) +{ + + int i,j,k; + for(i = 0; i < filters; ++i){ + mean_delta[i] = 0; + for (j = 0; j < batch; ++j) { + for (k = 0; k < spatial; ++k) { + int index = j*filters*spatial + i*spatial + k; + mean_delta[i] += delta[index]; + } + } + mean_delta[i] *= (-1./sqrt(variance[i] + .00001f)); + } +} +void variance_delta_cpu(float *x, float *delta, float *mean, float *variance, int batch, int filters, int spatial, float *variance_delta) +{ + + int i,j,k; + for(i = 0; i < filters; ++i){ + variance_delta[i] = 0; + for(j = 0; j < batch; ++j){ + for(k = 0; k < spatial; ++k){ + int index = j*filters*spatial + i*spatial + k; + variance_delta[i] += delta[index]*(x[index] - mean[i]); + } + } + variance_delta[i] *= -.5 * pow(variance[i] + .00001f, (float)(-3./2.)); + } +} +void normalize_delta_cpu(float *x, float *mean, float *variance, float *mean_delta, float *variance_delta, int batch, int filters, int spatial, float *delta) +{ + int f, j, k; + for(j = 0; j < batch; ++j){ + for(f = 0; f < filters; ++f){ + for(k = 0; k < spatial; ++k){ + int index = j*filters*spatial + f*spatial + k; + delta[index] = delta[index] * 1./(sqrt(variance[f]) + .00001f) + variance_delta[f] * 2. * (x[index] - mean[f]) / (spatial * batch) + mean_delta[f]/(spatial*batch); + } + } + } +} + +convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, int pad, ACTIVATION activation, int batch_normalize, int binary) { int i; convolutional_layer l = {0}; @@ -51,6 +109,7 @@ convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int l.w = w; l.c = c; l.n = n; + l.binary = binary; l.batch = batch; l.stride = stride; l.size = size; @@ -78,6 +137,10 @@ convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int l.output = calloc(l.batch*out_h * out_w * n, sizeof(float)); l.delta = calloc(l.batch*out_h * out_w * n, sizeof(float)); + if(binary){ + l.binary_filters = calloc(c*n*size*size, sizeof(float)); + } + if(batch_normalize){ l.scales = calloc(n, sizeof(float)); l.scale_updates = calloc(n, sizeof(float)); @@ -106,6 +169,10 @@ convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int l.delta_gpu = cuda_make_array(l.delta, l.batch*out_h*out_w*n); l.output_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n); + if(binary){ + l.binary_filters_gpu = cuda_make_array(l.filters, c*n*size*size); + } + if(batch_normalize){ l.mean_gpu = cuda_make_array(l.mean, n); l.variance_gpu = cuda_make_array(l.variance, n); @@ -141,7 +208,7 @@ void denormalize_convolutional_layer(convolutional_layer l) void test_convolutional_layer() { - convolutional_layer l = make_convolutional_layer(1, 5, 5, 3, 2, 5, 2, 1, LEAKY, 1); + convolutional_layer l = make_convolutional_layer(1, 5, 5, 3, 2, 5, 2, 1, LEAKY, 1, 0); l.batch_normalize = 1; float data[] = {1,1,1,1,1, 1,1,1,1,1, diff --git a/src/convolutional_layer.h b/src/convolutional_layer.h index e22a3960..7c391813 100644 --- a/src/convolutional_layer.h +++ b/src/convolutional_layer.h @@ -2,7 +2,6 @@ #define CONVOLUTIONAL_LAYER_H #include "cuda.h" -#include "params.h" #include "image.h" #include "activations.h" #include "layer.h" @@ -22,7 +21,7 @@ void add_bias_gpu(float *output, float *biases, int batch, int n, int size); void backward_bias_gpu(float *bias_updates, float *delta, int batch, int n, int size); #endif -convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, int pad, ACTIVATION activation, int batch_normalization); +convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, int pad, ACTIVATION activation, int batch_normalization, int binary); void denormalize_convolutional_layer(convolutional_layer l); void resize_convolutional_layer(convolutional_layer *layer, int w, int h); void forward_convolutional_layer(const convolutional_layer layer, network_state state); diff --git a/src/cost_layer.c b/src/cost_layer.c index 75934901..39ae8096 100644 --- a/src/cost_layer.c +++ b/src/cost_layer.c @@ -11,7 +11,8 @@ COST_TYPE get_cost_type(char *s) { if (strcmp(s, "sse")==0) return SSE; if (strcmp(s, "masked")==0) return MASKED; - fprintf(stderr, "Couldn't find activation function %s, going with SSE\n", s); + if (strcmp(s, "smooth")==0) return SMOOTH; + fprintf(stderr, "Couldn't find cost type %s, going with SSE\n", s); return SSE; } @@ -22,6 +23,8 @@ char *get_cost_string(COST_TYPE a) return "sse"; case MASKED: return "masked"; + case SMOOTH: + return "smooth"; } return "sse"; } @@ -65,8 +68,12 @@ void forward_cost_layer(cost_layer l, network_state state) if(state.truth[i] == SECRET_NUM) state.input[i] = SECRET_NUM; } } - copy_cpu(l.batch*l.inputs, state.truth, 1, l.delta, 1); - axpy_cpu(l.batch*l.inputs, -1, state.input, 1, l.delta, 1); + if(l.cost_type == SMOOTH){ + smooth_l1_cpu(l.batch*l.inputs, state.input, state.truth, l.delta); + } else { + copy_cpu(l.batch*l.inputs, state.truth, 1, l.delta, 1); + axpy_cpu(l.batch*l.inputs, -1, state.input, 1, l.delta, 1); + } *(l.output) = dot_cpu(l.batch*l.inputs, l.delta, 1, l.delta, 1); //printf("cost: %f\n", *l.output); } @@ -95,8 +102,12 @@ void forward_cost_layer_gpu(cost_layer l, network_state state) mask_ongpu(l.batch*l.inputs, state.input, SECRET_NUM, state.truth); } - copy_ongpu(l.batch*l.inputs, state.truth, 1, l.delta_gpu, 1); - axpy_ongpu(l.batch*l.inputs, -1, state.input, 1, l.delta_gpu, 1); + if(l.cost_type == SMOOTH){ + smooth_l1_gpu(l.batch*l.inputs, state.input, state.truth, l.delta_gpu); + } else { + copy_ongpu(l.batch*l.inputs, state.truth, 1, l.delta_gpu, 1); + axpy_ongpu(l.batch*l.inputs, -1, state.input, 1, l.delta_gpu, 1); + } cuda_pull_array(l.delta_gpu, l.delta, l.batch*l.inputs); *(l.output) = dot_cpu(l.batch*l.inputs, l.delta, 1, l.delta, 1); diff --git a/src/crop_layer.h b/src/crop_layer.h index 12112f02..3aa2d3dd 100644 --- a/src/crop_layer.h +++ b/src/crop_layer.h @@ -2,7 +2,6 @@ #define CROP_LAYER_H #include "image.h" -#include "params.h" #include "layer.h" #include "network.h" diff --git a/src/crop_layer_kernels.cu b/src/crop_layer_kernels.cu index da635a6f..8a086305 100644 --- a/src/crop_layer_kernels.cu +++ b/src/crop_layer_kernels.cu @@ -184,7 +184,7 @@ extern "C" void forward_crop_layer_gpu(crop_layer layer, network_state state) { cuda_random(layer.rand_gpu, layer.batch*8); - float radians = layer.angle*3.14159/180.; + float radians = layer.angle*3.14159265/180.; float scale = 2; float translate = -1; diff --git a/src/darknet.c b/src/darknet.c index c2a65965..938609ea 100644 --- a/src/darknet.c +++ b/src/darknet.c @@ -20,6 +20,7 @@ extern void run_nightmare(int argc, char **argv); extern void run_dice(int argc, char **argv); extern void run_compare(int argc, char **argv); extern void run_classifier(int argc, char **argv); +extern void run_char_rnn(int argc, char **argv); void change_rate(char *filename, float scale, float add) { @@ -203,7 +204,10 @@ int main(int argc, char **argv) return 0; } gpu_index = find_int_arg(argc, argv, "-i", 0); - if(find_arg(argc, argv, "-nogpu")) gpu_index = -1; + if(find_arg(argc, argv, "-nogpu")) { + gpu_index = -1; + printf("nogpu\n"); + } #ifndef GPU gpu_index = -1; @@ -220,6 +224,8 @@ int main(int argc, char **argv) average(argc, argv); } else if (0 == strcmp(argv[1], "yolo")){ run_yolo(argc, argv); + } else if (0 == strcmp(argv[1], "rnn")){ + run_char_rnn(argc, argv); } else if (0 == strcmp(argv[1], "coco")){ run_coco(argc, argv); } else if (0 == strcmp(argv[1], "classifier")){ diff --git a/src/deconvolutional_layer.h b/src/deconvolutional_layer.h index b6af3978..2d36e02a 100644 --- a/src/deconvolutional_layer.h +++ b/src/deconvolutional_layer.h @@ -2,7 +2,6 @@ #define DECONVOLUTIONAL_LAYER_H #include "cuda.h" -#include "params.h" #include "image.h" #include "activations.h" #include "layer.h" diff --git a/src/detection_layer.c b/src/detection_layer.c index ca32bc03..90b672b1 100644 --- a/src/detection_layer.c +++ b/src/detection_layer.c @@ -50,7 +50,7 @@ void forward_detection_layer(const detection_layer l, network_state state) int index = b*l.inputs; for (i = 0; i < locations; ++i) { int offset = i*l.classes; - softmax_array(l.output + index + offset, l.classes, + softmax_array(l.output + index + offset, l.classes, 1, l.output + index + offset); } int offset = locations*l.classes; diff --git a/src/dropout_layer.c b/src/dropout_layer.c index bb410dcf..29b9193c 100644 --- a/src/dropout_layer.c +++ b/src/dropout_layer.c @@ -1,5 +1,4 @@ #include "dropout_layer.h" -#include "params.h" #include "utils.h" #include "cuda.h" #include diff --git a/src/dropout_layer.h b/src/dropout_layer.h index 0c2ce4d8..691cfc5b 100644 --- a/src/dropout_layer.h +++ b/src/dropout_layer.h @@ -1,7 +1,6 @@ #ifndef DROPOUT_LAYER_H #define DROPOUT_LAYER_H -#include "params.h" #include "layer.h" #include "network.h" diff --git a/src/dropout_layer_kernels.cu b/src/dropout_layer_kernels.cu index f65e3c3d..7e51bd55 100644 --- a/src/dropout_layer_kernels.cu +++ b/src/dropout_layer_kernels.cu @@ -6,7 +6,6 @@ extern "C" { #include "dropout_layer.h" #include "cuda.h" #include "utils.h" -#include "params.h" } __global__ void yoloswag420blazeit360noscope(float *input, int size, float *rand, float prob, float scale) diff --git a/src/image.c b/src/image.c index d7d57d56..60ccfb8c 100644 --- a/src/image.c +++ b/src/image.c @@ -708,6 +708,8 @@ image resize_image(image im, int w, int h) void test_resize(char *filename) { image im = load_image(filename, 0,0, 3); + float mag = mag_array(im.data, im.w*im.h*im.c); + printf("L2 Norm: %f\n", mag); image gray = grayscale_image(im); image sat2 = copy_image(im); diff --git a/src/layer.h b/src/layer.h index d8af6e40..fc76234f 100644 --- a/src/layer.h +++ b/src/layer.h @@ -21,11 +21,12 @@ typedef enum { AVGPOOL, LOCAL, SHORTCUT, - ACTIVE + ACTIVE, + RNN } LAYER_TYPE; typedef enum{ - SSE, MASKED + SSE, MASKED, SMOOTH } COST_TYPE; struct layer{ @@ -50,6 +51,9 @@ struct layer{ int sqrt; int flip; int index; + int binary; + int steps; + int hidden; float angle; float jitter; float saturation; @@ -77,6 +81,7 @@ struct layer{ int dontload; int dontloadscales; + float temperature; float probability; float scale; @@ -85,6 +90,9 @@ struct layer{ float *cost; float *filters; float *filter_updates; + float *state; + + float *binary_filters; float *biases; float *bias_updates; @@ -107,14 +115,28 @@ struct layer{ float * mean; float * variance; + float * mean_delta; + float * variance_delta; + float * rolling_mean; float * rolling_variance; + float * x; + float * x_norm; + + struct layer *input_layer; + struct layer *self_layer; + struct layer *output_layer; + #ifdef GPU int *indexes_gpu; + float * state_gpu; float * filters_gpu; float * filter_updates_gpu; + float *binary_filters_gpu; + float *mean_filters_gpu; + float * spatial_mean_gpu; float * spatial_variance_gpu; diff --git a/src/maxpool_layer.h b/src/maxpool_layer.h index b91c2c18..4729e43a 100644 --- a/src/maxpool_layer.h +++ b/src/maxpool_layer.h @@ -2,7 +2,6 @@ #define MAXPOOL_LAYER_H #include "image.h" -#include "params.h" #include "cuda.h" #include "layer.h" #include "network.h" diff --git a/src/network.c b/src/network.c index 79579f10..32c3ba14 100644 --- a/src/network.c +++ b/src/network.c @@ -8,6 +8,7 @@ #include "crop_layer.h" #include "connected_layer.h" +#include "rnn_layer.h" #include "local_layer.h" #include "convolutional_layer.h" #include "activation_layer.h" @@ -82,6 +83,8 @@ char *get_layer_string(LAYER_TYPE a) return "deconvolutional"; case CONNECTED: return "connected"; + case RNN: + return "rnn"; case MAXPOOL: return "maxpool"; case AVGPOOL: @@ -144,6 +147,8 @@ void forward_network(network net, network_state state) forward_detection_layer(l, state); } else if(l.type == CONNECTED){ forward_connected_layer(l, state); + } else if(l.type == RNN){ + forward_rnn_layer(l, state); } else if(l.type == CROP){ forward_crop_layer(l, state); } else if(l.type == COST){ @@ -178,6 +183,8 @@ void update_network(network net) update_deconvolutional_layer(l, rate, net.momentum, net.decay); } else if(l.type == CONNECTED){ update_connected_layer(l, update_batch, rate, net.momentum, net.decay); + } else if(l.type == RNN){ + update_rnn_layer(l, update_batch, rate, net.momentum, net.decay); } else if(l.type == LOCAL){ update_local_layer(l, update_batch, rate, net.momentum, net.decay); } @@ -252,6 +259,8 @@ void backward_network(network net, network_state state) if(i != 0) backward_softmax_layer(l, state); } else if(l.type == CONNECTED){ backward_connected_layer(l, state); + } else if(l.type == RNN){ + backward_rnn_layer(l, state); } else if(l.type == LOCAL){ backward_local_layer(l, state); } else if(l.type == COST){ diff --git a/src/network.h b/src/network.h index 4c108df0..3d7c5746 100644 --- a/src/network.h +++ b/src/network.h @@ -5,7 +5,6 @@ #include "image.h" #include "layer.h" #include "data.h" -#include "params.h" typedef enum { CONSTANT, STEP, EXP, POLY, STEPS, SIG @@ -28,6 +27,7 @@ typedef struct network{ float gamma; float scale; float power; + int time_steps; int step; int max_batches; float *scales; @@ -77,6 +77,7 @@ void update_network(network net); float train_network(network net, data d); float train_network_batch(network net, data d, int n); float train_network_sgd(network net, data d, int n); +float train_network_datum(network net, float *x, float *y); matrix network_predict_data(network net, data test); float *network_predict(network net, float *input); diff --git a/src/network_kernels.cu b/src/network_kernels.cu index a83293da..ea128194 100644 --- a/src/network_kernels.cu +++ b/src/network_kernels.cu @@ -11,11 +11,11 @@ extern "C" { #include "image.h" #include "data.h" #include "utils.h" -#include "params.h" #include "parser.h" #include "crop_layer.h" #include "connected_layer.h" +#include "rnn_layer.h" #include "detection_layer.h" #include "convolutional_layer.h" #include "activation_layer.h" @@ -57,6 +57,8 @@ void forward_network_gpu(network net, network_state state) forward_detection_layer_gpu(l, state); } else if(l.type == CONNECTED){ forward_connected_layer_gpu(l, state); + } else if(l.type == RNN){ + forward_rnn_layer_gpu(l, state); } else if(l.type == CROP){ forward_crop_layer_gpu(l, state); } else if(l.type == COST){ @@ -118,6 +120,8 @@ void backward_network_gpu(network net, network_state state) if(i != 0) backward_softmax_layer_gpu(l, state); } else if(l.type == CONNECTED){ backward_connected_layer_gpu(l, state); + } else if(l.type == RNN){ + backward_rnn_layer_gpu(l, state); } else if(l.type == COST){ backward_cost_layer_gpu(l, state); } else if(l.type == ROUTE){ @@ -141,6 +145,8 @@ void update_network_gpu(network net) update_deconvolutional_layer_gpu(l, rate, net.momentum, net.decay); } else if(l.type == CONNECTED){ update_connected_layer_gpu(l, update_batch, rate, net.momentum, net.decay); + } else if(l.type == RNN){ + update_rnn_layer_gpu(l, update_batch, rate, net.momentum, net.decay); } else if(l.type == LOCAL){ update_local_layer_gpu(l, update_batch, rate, net.momentum, net.decay); } diff --git a/src/nightmare.c b/src/nightmare.c index ccc75255..bc2060f2 100644 --- a/src/nightmare.c +++ b/src/nightmare.c @@ -4,6 +4,10 @@ #include "blas.h" #include "utils.h" +#ifdef OPENCV +#include "opencv2/highgui/highgui_c.h" +#endif + float abs_mean(float *x, int n) { int i; @@ -167,6 +171,10 @@ void reconstruct_picture(network net, float *features, image recon, image update translate_image(recon, 1); scale_image(recon, .5); + + float mag = mag_array(recon.data, recon.w*recon.h*recon.c); + scal_cpu(recon.w*recon.h*recon.c, 600/mag, recon.data, 1); + constrain_image(recon); free_image(delta); } @@ -222,10 +230,21 @@ void run_nightmare(int argc, char **argv) image update; if (reconstruct){ resize_network(&net, im.w, im.h); - int size = get_network_output_size(net); - features = calloc(size, sizeof(float)); - float *out = network_predict(net, im.data); - copy_cpu(size, out, 1, features, 1); + + int zz = 0; + network_predict(net, im.data); + image out_im = get_network_image(net); + image crop = crop_image(out_im, zz, zz, out_im.w-2*zz, out_im.h-2*zz); + //flip_image(crop); + image f_im = resize_image(crop, out_im.w, out_im.h); + free_image(crop); + printf("%d features\n", out_im.w*out_im.h*out_im.c); + + + im = resize_image(im, im.w*2, im.h); + f_im = resize_image(f_im, f_im.w*2, f_im.h); + features = f_im.data; + free_image(im); im = make_random_image(im.w, im.h, im.c); update = make_image(im.w, im.h, im.c); diff --git a/src/params.h b/src/params.h deleted file mode 100644 index 8b137891..00000000 --- a/src/params.h +++ /dev/null @@ -1 +0,0 @@ - diff --git a/src/parser.c b/src/parser.c index 218fd27a..a48f207c 100644 --- a/src/parser.c +++ b/src/parser.c @@ -11,6 +11,7 @@ #include "normalization_layer.h" #include "deconvolutional_layer.h" #include "connected_layer.h" +#include "rnn_layer.h" #include "maxpool_layer.h" #include "softmax_layer.h" #include "dropout_layer.h" @@ -34,6 +35,7 @@ int is_activation(section *s); int is_local(section *s); int is_deconvolutional(section *s); int is_connected(section *s); +int is_rnn(section *s); int is_maxpool(section *s); int is_avgpool(section *s); int is_dropout(section *s); @@ -85,6 +87,7 @@ typedef struct size_params{ int w; int c; int index; + int time_steps; } size_params; deconvolutional_layer parse_deconvolutional(list *options, size_params params) @@ -151,8 +154,9 @@ convolutional_layer parse_convolutional(list *options, size_params params) batch=params.batch; if(!(h && w && c)) error("Layer before convolutional layer must output image."); int batch_normalize = option_find_int_quiet(options, "batch_normalize", 0); + int binary = option_find_int_quiet(options, "binary", 0); - convolutional_layer layer = make_convolutional_layer(batch,h,w,c,n,size,stride,pad,activation, batch_normalize); + convolutional_layer layer = make_convolutional_layer(batch,h,w,c,n,size,stride,pad,activation, batch_normalize, binary); layer.flipped = option_find_int_quiet(options, "flipped", 0); char *weights = option_find_str(options, "weights", 0); @@ -165,13 +169,27 @@ convolutional_layer parse_convolutional(list *options, size_params params) return layer; } +layer parse_rnn(list *options, size_params params) +{ + int output = option_find_int(options, "output",1); + int hidden = option_find_int(options, "hidden",1); + char *activation_s = option_find_str(options, "activation", "logistic"); + ACTIVATION activation = get_activation(activation_s); + int batch_normalize = option_find_int_quiet(options, "batch_normalize", 0); + + layer l = make_rnn_layer(params.batch, params.inputs, hidden, output, params.time_steps, activation, batch_normalize); + + return l; +} + connected_layer parse_connected(list *options, size_params params) { int output = option_find_int(options, "output",1); char *activation_s = option_find_str(options, "activation", "logistic"); ACTIVATION activation = get_activation(activation_s); + int batch_normalize = option_find_int_quiet(options, "batch_normalize", 0); - connected_layer layer = make_connected_layer(params.batch, params.inputs, output, activation); + connected_layer layer = make_connected_layer(params.batch, params.inputs, output, activation, batch_normalize); char *weights = option_find_str(options, "weights", 0); char *biases = option_find_str(options, "biases", 0); @@ -185,8 +203,9 @@ connected_layer parse_connected(list *options, size_params params) softmax_layer parse_softmax(list *options, size_params params) { - int groups = option_find_int(options, "groups",1); + int groups = option_find_int_quiet(options, "groups",1); softmax_layer layer = make_softmax_layer(params.batch, params.inputs, groups); + layer.temperature = option_find_float_quiet(options, "temperature", 1); return layer; } @@ -388,7 +407,9 @@ void parse_net_options(list *options, network *net) net->momentum = option_find_float(options, "momentum", .9); net->decay = option_find_float(options, "decay", .0001); int subdivs = option_find_int(options, "subdivisions",1); + net->time_steps = option_find_int_quiet(options, "time_steps",1); net->batch /= subdivs; + net->batch *= net->time_steps; net->subdivisions = subdivs; net->h = option_find_int_quiet(options, "height",0); @@ -456,6 +477,7 @@ network parse_network_cfg(char *filename) params.c = net.c; params.inputs = net.inputs; params.batch = net.batch; + params.time_steps = net.time_steps; n = n->next; int count = 0; @@ -474,6 +496,8 @@ network parse_network_cfg(char *filename) l = parse_activation(options, params); }else if(is_deconvolutional(s)){ l = parse_deconvolutional(options, params); + }else if(is_rnn(s)){ + l = parse_rnn(options, params); }else if(is_connected(s)){ l = parse_connected(options, params); }else if(is_crop(s)){ @@ -564,6 +588,10 @@ int is_network(section *s) return (strcmp(s->type, "[net]")==0 || strcmp(s->type, "[network]")==0); } +int is_rnn(section *s) +{ + return (strcmp(s->type, "[rnn]")==0); +} int is_connected(section *s) { return (strcmp(s->type, "[conn]")==0 @@ -674,6 +702,22 @@ void save_weights_double(network net, char *filename) fclose(fp); } +void save_connected_weights(layer l, FILE *fp) +{ +#ifdef GPU + if(gpu_index >= 0){ + pull_connected_layer(l); + } +#endif + fwrite(l.biases, sizeof(float), l.outputs, fp); + fwrite(l.weights, sizeof(float), l.outputs*l.inputs, fp); + if (l.batch_normalize){ + fwrite(l.scales, sizeof(float), l.outputs, fp); + fwrite(l.rolling_mean, sizeof(float), l.outputs, fp); + fwrite(l.rolling_variance, sizeof(float), l.outputs, fp); + } +} + void save_weights_upto(network net, char *filename, int cutoff) { fprintf(stderr, "Saving weights to %s\n", filename); @@ -706,13 +750,11 @@ void save_weights_upto(network net, char *filename, int cutoff) } fwrite(l.filters, sizeof(float), num, fp); } if(l.type == CONNECTED){ -#ifdef GPU - if(gpu_index >= 0){ - pull_connected_layer(l); - } -#endif - fwrite(l.biases, sizeof(float), l.outputs, fp); - fwrite(l.weights, sizeof(float), l.outputs*l.inputs, fp); + save_connected_weights(l, fp); + } if(l.type == RNN){ + save_connected_weights(*(l.input_layer), fp); + save_connected_weights(*(l.self_layer), fp); + save_connected_weights(*(l.output_layer), fp); } if(l.type == LOCAL){ #ifdef GPU if(gpu_index >= 0){ @@ -745,6 +787,25 @@ void transpose_matrix(float *a, int rows, int cols) free(transpose); } +void load_connected_weights(layer l, FILE *fp, int transpose) +{ + fread(l.biases, sizeof(float), l.outputs, fp); + fread(l.weights, sizeof(float), l.outputs*l.inputs, fp); + if(transpose){ + transpose_matrix(l.weights, l.inputs, l.outputs); + } + if (l.batch_normalize && (!l.dontloadscales)){ + fread(l.scales, sizeof(float), l.outputs, fp); + fread(l.rolling_mean, sizeof(float), l.outputs, fp); + fread(l.rolling_variance, sizeof(float), l.outputs, fp); + } +#ifdef GPU + if(gpu_index >= 0){ + push_connected_layer(l); + } +#endif +} + void load_weights_upto(network *net, char *filename, int cutoff) { fprintf(stderr, "Loading weights from %s...", filename); @@ -759,6 +820,7 @@ void load_weights_upto(network *net, char *filename, int cutoff) fread(&minor, sizeof(int), 1, fp); fread(&revision, sizeof(int), 1, fp); fread(net->seen, sizeof(int), 1, fp); + int transpose = (major > 1000) || (minor > 1000); int i; for(i = 0; i < net->n && i < cutoff; ++i){ @@ -793,16 +855,12 @@ void load_weights_upto(network *net, char *filename, int cutoff) #endif } if(l.type == CONNECTED){ - fread(l.biases, sizeof(float), l.outputs, fp); - fread(l.weights, sizeof(float), l.outputs*l.inputs, fp); - if(major > 1000 || minor > 1000){ - transpose_matrix(l.weights, l.inputs, l.outputs); - } -#ifdef GPU - if(gpu_index >= 0){ - push_connected_layer(l); - } -#endif + load_connected_weights(l, fp, transpose); + } + if(l.type == RNN){ + load_connected_weights(*(l.input_layer), fp, transpose); + load_connected_weights(*(l.self_layer), fp, transpose); + load_connected_weights(*(l.output_layer), fp, transpose); } if(l.type == LOCAL){ int locations = l.out_w*l.out_h; diff --git a/src/rnn.c b/src/rnn.c new file mode 100644 index 00000000..d3e7e51f --- /dev/null +++ b/src/rnn.c @@ -0,0 +1,147 @@ +#include "network.h" +#include "cost_layer.h" +#include "utils.h" +#include "parser.h" + +#ifdef OPENCV +#include "opencv2/highgui/highgui_c.h" +#endif + +typedef struct { + float *x; + float *y; +} float_pair; + +float_pair get_rnn_data(char *text, int len, int batch, int steps) +{ + float *x = calloc(batch * steps * 256, sizeof(float)); + float *y = calloc(batch * steps * 256, sizeof(float)); + int i,j; + for(i = 0; i < batch; ++i){ + int index = rand() %(len - steps - 1); + for(j = 0; j < steps; ++j){ + x[(j*batch + i)*256 + text[index + j]] = 1; + y[(j*batch + i)*256 + text[index + j + 1]] = 1; + } + } + float_pair p; + p.x = x; + p.y = y; + return p; +} + +void train_char_rnn(char *cfgfile, char *weightfile, char *filename) +{ + FILE *fp = fopen(filename, "r"); + //FILE *fp = fopen("data/ab.txt", "r"); + //FILE *fp = fopen("data/grrm/asoiaf.txt", "r"); + + fseek(fp, 0, SEEK_END); + size_t size = ftell(fp); + fseek(fp, 0, SEEK_SET); + + char *text = calloc(size, sizeof(char)); + fread(text, 1, size, fp); + fclose(fp); + + char *backup_directory = "/home/pjreddie/backup/"; + srand(time(0)); + data_seed = time(0); + char *base = basecfg(cfgfile); + printf("%s\n", base); + float avg_loss = -1; + network net = parse_network_cfg(cfgfile); + if(weightfile){ + load_weights(&net, weightfile); + } + printf("Learning Rate: %g, Momentum: %g, Decay: %g\n", net.learning_rate, net.momentum, net.decay); + int batch = net.batch; + int steps = net.time_steps; + int i = (*net.seen)/net.batch; + + clock_t time; + while(get_current_batch(net) < net.max_batches){ + i += 1; + time=clock(); + float_pair p = get_rnn_data(text, size, batch/steps, steps); + + float loss = train_network_datum(net, p.x, p.y) / (batch); + free(p.x); + free(p.y); + if (avg_loss < 0) avg_loss = loss; + avg_loss = avg_loss*.9 + loss*.1; + + printf("%d: %f, %f avg, %f rate, %lf seconds\n", i, loss, avg_loss, get_current_rate(net), sec(clock()-time)); + if(i%100==0){ + char buff[256]; + sprintf(buff, "%s/%s_%d.weights", backup_directory, base, i); + save_weights(net, buff); + } + if(i%10==0){ + char buff[256]; + sprintf(buff, "%s/%s.backup", backup_directory, base); + save_weights(net, buff); + } + } + char buff[256]; + sprintf(buff, "%s/%s_final.weights", backup_directory, base); + save_weights(net, buff); +} + +void test_char_rnn(char *cfgfile, char *weightfile, int num, char *seed, float temp, int rseed) +{ + srand(rseed); + char *base = basecfg(cfgfile); + printf("%s\n", base); + + network net = parse_network_cfg(cfgfile); + if(weightfile){ + load_weights(&net, weightfile); + } + + int i, j; + for(i = 0; i < net.n; ++i) net.layers[i].temperature = temp; + char c; + int len = strlen(seed); + float *input = calloc(256, sizeof(float)); + for(i = 0; i < len-1; ++i){ + c = seed[i]; + input[(int)c] = 1; + network_predict(net, input); + input[(int)c] = 0; + printf("%c", c); + } + c = seed[len-1]; + for(i = 0; i < num; ++i){ + printf("%c", c); + float r = rand_uniform(0,1); + float sum = 0; + input[(int)c] = 1; + float *out = network_predict(net, input); + input[(int)c] = 0; + for(j = 0; j < 256; ++j){ + sum += out[j]; + if(sum > r) break; + } + c = j; + } + printf("\n"); +} + +void run_char_rnn(int argc, char **argv) +{ + if(argc < 4){ + fprintf(stderr, "usage: %s %s [train/test/valid] [cfg] [weights (optional)]\n", argv[0], argv[1]); + return; + } + char *filename = find_char_arg(argc, argv, "-file", "data/shakespeare.txt"); + char *seed = find_char_arg(argc, argv, "-seed", "\n"); + int len = find_int_arg(argc, argv, "-len", 100); + float temp = find_float_arg(argc, argv, "-temp", 1); + int rseed = find_int_arg(argc, argv, "-srand", time(0)); + + char *cfg = argv[3]; + char *weights = (argc > 4) ? argv[4] : 0; + if(0==strcmp(argv[2], "train")) train_char_rnn(cfg, weights, filename); + else if(0==strcmp(argv[2], "test")) test_char_rnn(cfg, weights, len, seed, temp, rseed); +} diff --git a/src/rnn_layer.c b/src/rnn_layer.c new file mode 100644 index 00000000..63582858 --- /dev/null +++ b/src/rnn_layer.c @@ -0,0 +1,275 @@ +#include "rnn_layer.h" +#include "connected_layer.h" +#include "utils.h" +#include "cuda.h" +#include "blas.h" +#include "gemm.h" + +#include +#include +#include +#include + + +layer make_rnn_layer(int batch, int inputs, int hidden, int outputs, int steps, ACTIVATION activation, int batch_normalize) +{ + printf("%d %d\n", batch, steps); + batch = batch / steps; + layer l = {0}; + l.batch = batch; + l.type = RNN; + l.steps = steps; + l.hidden = hidden; + l.inputs = inputs; + + l.state = calloc(batch*hidden, sizeof(float)); + + l.input_layer = malloc(sizeof(layer)); + *(l.input_layer) = make_connected_layer(batch*steps, inputs, hidden, activation, batch_normalize); + l.input_layer->batch = batch; + + l.self_layer = malloc(sizeof(layer)); + *(l.self_layer) = make_connected_layer(batch*steps, hidden, hidden, activation, batch_normalize); + l.self_layer->batch = batch; + + l.output_layer = malloc(sizeof(layer)); + *(l.output_layer) = make_connected_layer(batch*steps, hidden, outputs, activation, batch_normalize); + l.output_layer->batch = batch; + + l.outputs = outputs; + l.output = l.output_layer->output; + l.delta = l.output_layer->delta; + + #ifdef GPU + l.state_gpu = cuda_make_array(l.state, batch*hidden); + l.output_gpu = l.output_layer->output_gpu; + l.delta_gpu = l.output_layer->delta_gpu; + #endif + + fprintf(stderr, "RNN Layer: %d inputs, %d outputs\n", inputs, outputs); + return l; +} + +void update_rnn_layer(layer l, int batch, float learning_rate, float momentum, float decay) +{ + update_connected_layer(*(l.input_layer), batch, learning_rate, momentum, decay); + update_connected_layer(*(l.self_layer), batch, learning_rate, momentum, decay); + update_connected_layer(*(l.output_layer), batch, learning_rate, momentum, decay); +} + +void forward_rnn_layer(layer l, network_state state) +{ + network_state s = {0}; + s.train = state.train; + int i; + layer input_layer = *(l.input_layer); + layer self_layer = *(l.self_layer); + layer output_layer = *(l.output_layer); + + fill_cpu(l.outputs * l.batch * l.steps, 0, output_layer.delta, 1); + fill_cpu(l.hidden * l.batch * l.steps, 0, self_layer.delta, 1); + fill_cpu(l.hidden * l.batch * l.steps, 0, input_layer.delta, 1); + if(state.train) fill_cpu(l.hidden * l.batch, 0, l.state, 1); + + for (i = 0; i < l.steps; ++i) { + s.input = state.input; + forward_connected_layer(input_layer, s); + + s.input = l.state; + forward_connected_layer(self_layer, s); + + copy_cpu(l.hidden * l.batch, input_layer.output, 1, l.state, 1); + axpy_cpu(l.hidden * l.batch, 1, self_layer.output, 1, l.state, 1); + + s.input = l.state; + forward_connected_layer(output_layer, s); + + state.input += l.inputs*l.batch; + input_layer.output += l.hidden*l.batch; + self_layer.output += l.hidden*l.batch; + output_layer.output += l.outputs*l.batch; + } +} + +void backward_rnn_layer(layer l, network_state state) +{ + network_state s = {0}; + s.train = state.train; + int i; + layer input_layer = *(l.input_layer); + layer self_layer = *(l.self_layer); + layer output_layer = *(l.output_layer); + input_layer.output += l.hidden*l.batch*(l.steps-1); + input_layer.delta += l.hidden*l.batch*(l.steps-1); + + self_layer.output += l.hidden*l.batch*(l.steps-1); + self_layer.delta += l.hidden*l.batch*(l.steps-1); + + output_layer.output += l.outputs*l.batch*(l.steps-1); + output_layer.delta += l.outputs*l.batch*(l.steps-1); + for (i = l.steps-1; i >= 0; --i) { + copy_cpu(l.hidden * l.batch, input_layer.output, 1, l.state, 1); + axpy_cpu(l.hidden * l.batch, 1, self_layer.output, 1, l.state, 1); + + s.input = l.state; + s.delta = self_layer.delta; + backward_connected_layer(output_layer, s); + + if(i > 0){ + copy_cpu(l.hidden * l.batch, input_layer.output - l.hidden*l.batch, 1, l.state, 1); + axpy_cpu(l.hidden * l.batch, 1, self_layer.output - l.hidden*l.batch, 1, l.state, 1); + }else{ + fill_cpu(l.hidden * l.batch, 0, l.state, 1); + } + + s.input = l.state; + s.delta = self_layer.delta - l.hidden*l.batch; + if (i == 0) s.delta = 0; + backward_connected_layer(self_layer, s); + + copy_cpu(l.hidden*l.batch, self_layer.delta, 1, input_layer.delta, 1); + s.input = state.input + i*l.inputs*l.batch; + if(state.delta) s.delta = state.delta + i*l.inputs*l.batch; + else s.delta = 0; + backward_connected_layer(input_layer, s); + + input_layer.output -= l.hidden*l.batch; + input_layer.delta -= l.hidden*l.batch; + + self_layer.output -= l.hidden*l.batch; + self_layer.delta -= l.hidden*l.batch; + + output_layer.output -= l.outputs*l.batch; + output_layer.delta -= l.outputs*l.batch; + } +} + +#ifdef GPU + +void pull_rnn_layer(layer l) +{ + pull_connected_layer(*(l.input_layer)); + pull_connected_layer(*(l.self_layer)); + pull_connected_layer(*(l.output_layer)); +} + +void push_rnn_layer(layer l) +{ + push_connected_layer(*(l.input_layer)); + push_connected_layer(*(l.self_layer)); + push_connected_layer(*(l.output_layer)); +} + +void update_rnn_layer_gpu(layer l, int batch, float learning_rate, float momentum, float decay) +{ + update_connected_layer_gpu(*(l.input_layer), batch, learning_rate, momentum, decay); + update_connected_layer_gpu(*(l.self_layer), batch, learning_rate, momentum, decay); + update_connected_layer_gpu(*(l.output_layer), batch, learning_rate, momentum, decay); +} + +void forward_rnn_layer_gpu(layer l, network_state state) +{ + network_state s = {0}; + s.train = state.train; + int i; + layer input_layer = *(l.input_layer); + layer self_layer = *(l.self_layer); + layer output_layer = *(l.output_layer); + + fill_ongpu(l.outputs * l.batch * l.steps, 0, output_layer.delta_gpu, 1); + fill_ongpu(l.hidden * l.batch * l.steps, 0, self_layer.delta_gpu, 1); + fill_ongpu(l.hidden * l.batch * l.steps, 0, input_layer.delta_gpu, 1); + if(state.train) fill_ongpu(l.hidden * l.batch, 0, l.state_gpu, 1); + + for (i = 0; i < l.steps; ++i) { + s.input = state.input; + forward_connected_layer_gpu(input_layer, s); + + s.input = l.state_gpu; + forward_connected_layer_gpu(self_layer, s); + + copy_ongpu(l.hidden * l.batch, input_layer.output_gpu, 1, l.state_gpu, 1); + axpy_ongpu(l.hidden * l.batch, 1, self_layer.output_gpu, 1, l.state_gpu, 1); + + forward_connected_layer_gpu(output_layer, s); + + state.input += l.inputs*l.batch; + input_layer.output_gpu += l.hidden*l.batch; + input_layer.x_gpu += l.hidden*l.batch; + input_layer.x_norm_gpu += l.hidden*l.batch; + + self_layer.output_gpu += l.hidden*l.batch; + self_layer.x_gpu += l.hidden*l.batch; + self_layer.x_norm_gpu += l.hidden*l.batch; + + output_layer.output_gpu += l.outputs*l.batch; + output_layer.x_gpu += l.outputs*l.batch; + output_layer.x_norm_gpu += l.outputs*l.batch; + } +} + +void backward_rnn_layer_gpu(layer l, network_state state) +{ + network_state s = {0}; + s.train = state.train; + int i; + layer input_layer = *(l.input_layer); + layer self_layer = *(l.self_layer); + layer output_layer = *(l.output_layer); + input_layer.output_gpu += l.hidden*l.batch*(l.steps-1); + input_layer.delta_gpu += l.hidden*l.batch*(l.steps-1); + input_layer.x_gpu += l.hidden*l.batch*(l.steps-1); + input_layer.x_norm_gpu += l.hidden*l.batch*(l.steps-1); + + self_layer.output_gpu += l.hidden*l.batch*(l.steps-1); + self_layer.delta_gpu += l.hidden*l.batch*(l.steps-1); + self_layer.x_gpu += l.hidden*l.batch*(l.steps-1); + self_layer.x_norm_gpu += l.hidden*l.batch*(l.steps-1); + + output_layer.output_gpu += l.outputs*l.batch*(l.steps-1); + output_layer.delta_gpu += l.outputs*l.batch*(l.steps-1); + output_layer.x_gpu += l.outputs*l.batch*(l.steps-1); + output_layer.x_norm_gpu += l.outputs*l.batch*(l.steps-1); + for (i = l.steps-1; i >= 0; --i) { + copy_ongpu(l.hidden * l.batch, input_layer.output_gpu, 1, l.state_gpu, 1); + axpy_ongpu(l.hidden * l.batch, 1, self_layer.output_gpu, 1, l.state_gpu, 1); + + s.input = l.state_gpu; + s.delta = self_layer.delta_gpu; + backward_connected_layer_gpu(output_layer, s); + + if(i > 0){ + copy_ongpu(l.hidden * l.batch, input_layer.output_gpu - l.hidden*l.batch, 1, l.state_gpu, 1); + axpy_ongpu(l.hidden * l.batch, 1, self_layer.output_gpu - l.hidden*l.batch, 1, l.state_gpu, 1); + }else{ + fill_ongpu(l.hidden * l.batch, 0, l.state_gpu, 1); + } + + s.input = l.state_gpu; + s.delta = self_layer.delta_gpu - l.hidden*l.batch; + if (i == 0) s.delta = 0; + backward_connected_layer_gpu(self_layer, s); + + copy_ongpu(l.hidden*l.batch, self_layer.delta_gpu, 1, input_layer.delta_gpu, 1); + s.input = state.input + i*l.inputs*l.batch; + if(state.delta) s.delta = state.delta + i*l.inputs*l.batch; + else s.delta = 0; + backward_connected_layer_gpu(input_layer, s); + + input_layer.output_gpu -= l.hidden*l.batch; + input_layer.delta_gpu -= l.hidden*l.batch; + input_layer.x_gpu -= l.hidden*l.batch; + input_layer.x_norm_gpu -= l.hidden*l.batch; + + self_layer.output_gpu -= l.hidden*l.batch; + self_layer.delta_gpu -= l.hidden*l.batch; + self_layer.x_gpu -= l.hidden*l.batch; + self_layer.x_norm_gpu -= l.hidden*l.batch; + + output_layer.output_gpu -= l.outputs*l.batch; + output_layer.delta_gpu -= l.outputs*l.batch; + output_layer.x_gpu -= l.outputs*l.batch; + output_layer.x_norm_gpu -= l.outputs*l.batch; + } +} +#endif diff --git a/src/rnn_layer.h b/src/rnn_layer.h new file mode 100644 index 00000000..8d4f5854 --- /dev/null +++ b/src/rnn_layer.h @@ -0,0 +1,24 @@ + +#ifndef RNN_LAYER_H +#define RNN_LAYER_H + +#include "activations.h" +#include "layer.h" +#include "network.h" + +layer make_rnn_layer(int batch, int inputs, int hidden, int outputs, int steps, ACTIVATION activation, int batch_normalize); + +void forward_rnn_layer(layer l, network_state state); +void backward_rnn_layer(layer l, network_state state); +void update_rnn_layer(layer l, int batch, float learning_rate, float momentum, float decay); + +#ifdef GPU +void forward_rnn_layer_gpu(layer l, network_state state); +void backward_rnn_layer_gpu(layer l, network_state state); +void update_rnn_layer_gpu(layer l, int batch, float learning_rate, float momentum, float decay); +void push_rnn_layer(layer l); +void pull_rnn_layer(layer l); +#endif + +#endif + diff --git a/src/softmax_layer.c b/src/softmax_layer.c index 0d19acad..e189701f 100644 --- a/src/softmax_layer.c +++ b/src/softmax_layer.c @@ -26,7 +26,7 @@ softmax_layer make_softmax_layer(int batch, int inputs, int groups) return l; } -void softmax_array(float *input, int n, float *output) +void softmax_array(float *input, int n, float temp, float *output) { int i; float sum = 0; @@ -35,12 +35,12 @@ void softmax_array(float *input, int n, float *output) if(input[i] > largest) largest = input[i]; } for(i = 0; i < n; ++i){ - sum += exp(input[i]-largest); + sum += exp(input[i]/temp-largest/temp); } - if(sum) sum = largest+log(sum); + if(sum) sum = largest/temp+log(sum); else sum = largest-100; for(i = 0; i < n; ++i){ - output[i] = exp(input[i]-sum); + output[i] = exp(input[i]/temp-sum); } } @@ -50,7 +50,7 @@ void forward_softmax_layer(const softmax_layer l, network_state state) int inputs = l.inputs / l.groups; int batch = l.batch * l.groups; for(b = 0; b < batch; ++b){ - softmax_array(state.input+b*inputs, inputs, l.output+b*inputs); + softmax_array(state.input+b*inputs, inputs, l.temperature, l.output+b*inputs); } } diff --git a/src/softmax_layer.h b/src/softmax_layer.h index 9cbcd699..821a8dd7 100644 --- a/src/softmax_layer.h +++ b/src/softmax_layer.h @@ -1,12 +1,11 @@ #ifndef SOFTMAX_LAYER_H #define SOFTMAX_LAYER_H -#include "params.h" #include "layer.h" #include "network.h" typedef layer softmax_layer; -void softmax_array(float *input, int n, float *output); +void softmax_array(float *input, int n, float temp, float *output); softmax_layer make_softmax_layer(int batch, int inputs, int groups); void forward_softmax_layer(const softmax_layer l, network_state state); void backward_softmax_layer(const softmax_layer l, network_state state); diff --git a/src/softmax_layer_kernels.cu b/src/softmax_layer_kernels.cu index 66371805..8feaf89b 100644 --- a/src/softmax_layer_kernels.cu +++ b/src/softmax_layer_kernels.cu @@ -8,7 +8,7 @@ extern "C" { #include "blas.h" } -__global__ void forward_softmax_layer_kernel(int n, int batch, float *input, float *output) +__global__ void forward_softmax_layer_kernel(int n, int batch, float *input, float temp, float *output) { int b = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; if(b >= batch) return; @@ -21,11 +21,11 @@ __global__ void forward_softmax_layer_kernel(int n, int batch, float *input, flo largest = (val>largest) ? val : largest; } for(i = 0; i < n; ++i){ - sum += exp(input[i+b*n]-largest); + sum += exp(input[i+b*n]/temp-largest/temp); } - sum = (sum != 0) ? largest+log(sum) : largest-100; + sum = (sum != 0) ? largest/temp+log(sum) : largest-100; for(i = 0; i < n; ++i){ - output[i+b*n] = exp(input[i+b*n]-sum); + output[i+b*n] = exp(input[i+b*n]/temp-sum); } } @@ -38,7 +38,7 @@ extern "C" void forward_softmax_layer_gpu(const softmax_layer layer, network_sta { int inputs = layer.inputs / layer.groups; int batch = layer.batch * layer.groups; - forward_softmax_layer_kernel<<>>(inputs, batch, state.input, layer.output_gpu); + forward_softmax_layer_kernel<<>>(inputs, batch, state.input, layer.temperature, layer.output_gpu); check_error(cudaPeekAtLastError()); } diff --git a/src/utils.c b/src/utils.c index d49d0ce9..ec87a265 100644 --- a/src/utils.c +++ b/src/utils.c @@ -127,14 +127,13 @@ void pm(int M, int N, float *A) for(i =0 ; i < M; ++i){ printf("%d ", i+1); for(j = 0; j < N; ++j){ - printf("%10.6f, ", A[i*N+j]); + printf("%2.4f, ", A[i*N+j]); } printf("\n"); } printf("\n"); } - char *find_replace(char *str, char *orig, char *rep) { static char buffer[4096]; diff --git a/src/yolo.c b/src/yolo.c index 6bd4e6b2..382cbaa9 100644 --- a/src/yolo.c +++ b/src/yolo.c @@ -343,8 +343,10 @@ void test_yolo(char *cfgfile, char *weightfile, char *filename, float thresh) printf("%s: Predicted in %f seconds.\n", input, sec(clock()-time)); convert_yolo_detections(predictions, l.classes, l.n, l.sqrt, l.side, 1, 1, thresh, probs, boxes, 0); if (nms) do_nms_sort(boxes, probs, l.side*l.side*l.n, l.classes, nms); - draw_detections(im, l.side*l.side*l.n, thresh, boxes, probs, voc_names, voc_labels, 20); + //draw_detections(im, l.side*l.side*l.n, thresh, boxes, probs, voc_names, voc_labels, 20); + draw_detections(im, l.side*l.side*l.n, thresh, boxes, probs, voc_names, 0, 20); show_image(im, "predictions"); + save_image(im, "predictions"); show_image(sized, "resized"); free_image(im);