lots of stuff

Joseph Redmon

2016-01-28 913d355ec1cf34aad71fdd75202fc3b0309e63a0

lots of stuff

 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

 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;

 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

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

 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;
    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,6 +235,28 @@
    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);

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

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

 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,

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

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

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

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


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

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

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

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

 src/params.h

File was deleted

 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);
            load_connected_weights(l, fp, transpose);
            }
#ifdef GPU
            if(gpu_index >= 0){
                push_connected_layer(l);
            }
#endif
        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

New file
@@ -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

New file
@@ -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

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


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


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

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