From 913d355ec1cf34aad71fdd75202fc3b0309e63a0 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Thu, 28 Jan 2016 20:30:38 +0000
Subject: [PATCH] lots of stuff
---
src/dropout_layer_kernels.cu | 1
src/yolo.c | 4
src/softmax_layer_kernels.cu | 10
Makefile | 2
src/rnn.c | 147 ++++++++
src/nightmare.c | 27 +
src/convolutional_layer.h | 3
src/image.c | 2
src/crop_layer.h | 1
src/deconvolutional_layer.h | 1
src/blas.c | 12
src/convolutional_kernels.cu | 42 ++
src/dropout_layer.c | 1
src/layer.h | 26 +
src/utils.c | 3
src/network.c | 9
src/maxpool_layer.h | 1
src/cost_layer.c | 21
src/softmax_layer.h | 3
src/network.h | 3
src/dropout_layer.h | 1
src/network_kernels.cu | 8
src/connected_layer.c | 151 ++++++++
src/connected_layer.h | 2
src/rnn_layer.h | 24 +
src/softmax_layer.c | 10
src/blas.h | 11
src/detection_layer.c | 2
/dev/null | 1
src/rnn_layer.c | 275 ++++++++++++++++
src/convolutional_layer.c | 71 ++++
src/parser.c | 98 ++++-
src/crop_layer_kernels.cu | 2
src/blas_kernels.cu | 16
src/darknet.c | 8
35 files changed, 913 insertions(+), 86 deletions(-)
diff --git a/Makefile b/Makefile
index 04e30f3..c9b6eca 100644
--- a/Makefile
+++ b/Makefile
@@ -34,7 +34,7 @@
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 8769df3..d7948bb 100644
--- a/src/blas.c
+++ b/src/blas.c
@@ -67,7 +67,7 @@
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 @@
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 aecdc59..f5189e5 100644
--- a/src/blas.h
+++ b/src/blas.h
@@ -17,11 +17,18 @@
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_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 49406db..61db29f 100644
--- a/src/blas_kernels.cu
+++ b/src/blas_kernels.cu
@@ -409,3 +409,19 @@
shortcut_kernel<<<cuda_gridsize(size), BLOCK>>>(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<<<cuda_gridsize(n), BLOCK>>>(n, pred, truth, delta);
+ check_error(cudaPeekAtLastError());
+}
diff --git a/src/connected_layer.c b/src/connected_layer.c
index c0a9d8b..df78e67 100644
--- a/src/connected_layer.c
+++ b/src/connected_layer.c
@@ -9,7 +9,7 @@
#include <stdlib.h>
#include <string.h>
-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 @@
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 @@
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 @@
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 @@
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 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 @@
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 @@
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 @@
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 @@
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 @@
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 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 @@
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 @@
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 2bf53b2..56bd1c3 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 a64a499..4fdc1a1 100644
--- a/src/convolutional_kernels.cu
+++ b/src/convolutional_kernels.cu
@@ -12,6 +12,21 @@
#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 @@
}
}
+void binarize_filters_gpu(float *filters, int n, int size, float *mean)
+{
+ binarize_filters_kernel<<<cuda_gridsize(n), BLOCK>>>(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<<<n, BLOCK>>>(x_norm, delta, batch, n, size, scale_updates);
@@ -100,6 +121,13 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 871a84e..159951d 100644
--- a/src/convolutional_layer.c
+++ b/src/convolutional_layer.c
@@ -41,7 +41,65 @@
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 @@
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 @@
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 @@
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 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 e22a396..7c39181 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 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 7593490..39ae809 100644
--- a/src/cost_layer.c
+++ b/src/cost_layer.c
@@ -11,7 +11,8 @@
{
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 @@
return "sse";
case MASKED:
return "masked";
+ case SMOOTH:
+ return "smooth";
}
return "sse";
}
@@ -65,8 +68,12 @@
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 @@
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 12112f0..3aa2d3d 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 da635a6..8a08630 100644
--- a/src/crop_layer_kernels.cu
+++ b/src/crop_layer_kernels.cu
@@ -184,7 +184,7 @@
{
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 c2a6596..938609e 100644
--- a/src/darknet.c
+++ b/src/darknet.c
@@ -20,6 +20,7 @@
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 @@
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 @@
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 b6af397..2d36e02 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 ca32bc0..90b672b 100644
--- a/src/detection_layer.c
+++ b/src/detection_layer.c
@@ -50,7 +50,7 @@
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 bb410dc..29b9193 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 <stdlib.h>
diff --git a/src/dropout_layer.h b/src/dropout_layer.h
index 0c2ce4d..691cfc5 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 f65e3c3..7e51bd5 100644
--- a/src/dropout_layer_kernels.cu
+++ b/src/dropout_layer_kernels.cu
@@ -6,7 +6,6 @@
#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 d7d57d5..60ccfb8 100644
--- a/src/image.c
+++ b/src/image.c
@@ -708,6 +708,8 @@
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 d8af6e4..fc76234 100644
--- a/src/layer.h
+++ b/src/layer.h
@@ -21,11 +21,12 @@
AVGPOOL,
LOCAL,
SHORTCUT,
- ACTIVE
+ ACTIVE,
+ RNN
} LAYER_TYPE;
typedef enum{
- SSE, MASKED
+ SSE, MASKED, SMOOTH
} COST_TYPE;
struct layer{
@@ -50,6 +51,9 @@
int sqrt;
int flip;
int index;
+ int binary;
+ int steps;
+ int hidden;
float angle;
float jitter;
float saturation;
@@ -77,6 +81,7 @@
int dontload;
int dontloadscales;
+ float temperature;
float probability;
float scale;
@@ -85,6 +90,9 @@
float *cost;
float *filters;
float *filter_updates;
+ float *state;
+
+ float *binary_filters;
float *biases;
float *bias_updates;
@@ -107,14 +115,28 @@
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 b91c2c1..4729e43 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 79579f1..32c3ba1 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 @@
return "deconvolutional";
case CONNECTED:
return "connected";
+ case RNN:
+ return "rnn";
case MAXPOOL:
return "maxpool";
case AVGPOOL:
@@ -144,6 +147,8 @@
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 @@
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 @@
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 4c108df..3d7c574 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 @@
float gamma;
float scale;
float power;
+ int time_steps;
int step;
int max_batches;
float *scales;
@@ -77,6 +77,7 @@
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 a83293d..ea12819 100644
--- a/src/network_kernels.cu
+++ b/src/network_kernels.cu
@@ -11,11 +11,11 @@
#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 @@
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 @@
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 @@
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 ccc7525..bc2060f 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 @@
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 @@
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 8b13789..0000000
--- a/src/params.h
+++ /dev/null
@@ -1 +0,0 @@
-
diff --git a/src/parser.c b/src/parser.c
index 218fd27..a48f207 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_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 @@
int w;
int c;
int index;
+ int time_steps;
} size_params;
deconvolutional_layer parse_deconvolutional(list *options, size_params params)
@@ -151,8 +154,9 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
}
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 @@
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 @@
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 @@
#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 0000000..d3e7e51
--- /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 0000000..6358285
--- /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 <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+
+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 0000000..8d4f585
--- /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 0d19aca..e189701 100644
--- a/src/softmax_layer.c
+++ b/src/softmax_layer.c
@@ -26,7 +26,7 @@
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 @@
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 @@
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 9cbcd69..821a8dd 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 6637180..8feaf89 100644
--- a/src/softmax_layer_kernels.cu
+++ b/src/softmax_layer_kernels.cu
@@ -8,7 +8,7 @@
#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 @@
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 @@
{
int inputs = layer.inputs / layer.groups;
int batch = layer.batch * layer.groups;
- forward_softmax_layer_kernel<<<cuda_gridsize(batch), BLOCK>>>(inputs, batch, state.input, layer.output_gpu);
+ forward_softmax_layer_kernel<<<cuda_gridsize(batch), BLOCK>>>(inputs, batch, state.input, layer.temperature, layer.output_gpu);
check_error(cudaPeekAtLastError());
}
diff --git a/src/utils.c b/src/utils.c
index d49d0ce..ec87a26 100644
--- a/src/utils.c
+++ b/src/utils.c
@@ -127,14 +127,13 @@
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 6bd4e6b..382cbaa 100644
--- a/src/yolo.c
+++ b/src/yolo.c
@@ -343,8 +343,10 @@
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);
--
Gitblit v1.10.0