lots of stuff

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

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;

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

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<<<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());
+}

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 @@ 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;

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);

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<<<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 @@ 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);
         }
     }
 }

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,

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);

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);

src/crop_layer.h

@@ -2,7 +2,6 @@
 #define CROP_LAYER_H
 
 #include "image.h"
-#include "params.h"
 #include "layer.h"
 #include "network.h"
 

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;

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")){

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"

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;

src/dropout_layer.c

@@ -1,5 +1,4 @@
 #include "dropout_layer.h"
-#include "params.h"
 #include "utils.h"
 #include "cuda.h"
 #include <stdlib.h>

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"
 

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)

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);

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;
 

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"

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){

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);

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);
         }

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);

src/params.h

@@ -1 +0,0 @@
-

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;

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);
+}

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

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
+

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);
     }
 }
 

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);

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<<<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());
 }
 

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];

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);