Joseph Redmon
2015-03-23 7100de0b59892351c03ce60760431b24947e2ca3
going to break stuff
5 files modified
55 ■■■■■ changed files
src/convolutional_kernels.cu 14 ●●●● patch | view | raw | blame | history
src/detection.c 13 ●●●● patch | view | raw | blame | history
src/dropout_layer_kernels.cu 4 ●●●● patch | view | raw | blame | history
src/network.c 18 ●●●●● patch | view | raw | blame | history
src/network_kernels.cu 6 ●●●● patch | view | raw | blame | history
src/convolutional_kernels.cu
@@ -17,7 +17,7 @@
    if(offset < size) output[(batch*n+filter)*size + offset] = biases[filter];
}
extern "C" void bias_output_gpu(float *output, float *biases, int batch, int n, int size)
void bias_output_gpu(float *output, float *biases, int batch, int n, int size)
{
    dim3 dimBlock(BLOCK, 1, 1);
    dim3 dimGrid((size-1)/BLOCK + 1, n, batch);
@@ -46,13 +46,13 @@
    }
}
extern "C" void backward_bias_gpu(float *bias_updates, float *delta, int batch, int n, int size)
void backward_bias_gpu(float *bias_updates, float *delta, int batch, int n, int size)
{
    backward_bias_kernel<<<n, BLOCK>>>(bias_updates, delta, batch, n, size, 1);
    check_error(cudaPeekAtLastError());
}
extern "C" void forward_convolutional_layer_gpu(convolutional_layer layer, network_state state)
void forward_convolutional_layer_gpu(convolutional_layer layer, network_state state)
{
    int i;
    int m = layer.n;
@@ -71,7 +71,7 @@
    activate_array_ongpu(layer.output_gpu, m*n*layer.batch, layer.activation);
}
extern "C" void backward_convolutional_layer_gpu(convolutional_layer layer, network_state state)
void backward_convolutional_layer_gpu(convolutional_layer layer, network_state state)
{
    int i;
    int m = layer.n;
@@ -105,7 +105,7 @@
    }
}
extern "C" void pull_convolutional_layer(convolutional_layer layer)
void pull_convolutional_layer(convolutional_layer layer)
{
    cuda_pull_array(layer.filters_gpu, layer.filters, layer.c*layer.n*layer.size*layer.size);
    cuda_pull_array(layer.biases_gpu, layer.biases, layer.n);
@@ -113,7 +113,7 @@
    cuda_pull_array(layer.bias_updates_gpu, layer.bias_updates, layer.n);
}
extern "C" void push_convolutional_layer(convolutional_layer layer)
void push_convolutional_layer(convolutional_layer layer)
{
    cuda_push_array(layer.filters_gpu, layer.filters, layer.c*layer.n*layer.size*layer.size);
    cuda_push_array(layer.biases_gpu, layer.biases, layer.n);
@@ -121,7 +121,7 @@
    cuda_push_array(layer.bias_updates_gpu, layer.bias_updates, layer.n);
}
extern "C" void update_convolutional_layer_gpu(convolutional_layer layer, int batch, float learning_rate, float momentum, float decay)
void update_convolutional_layer_gpu(convolutional_layer layer, int batch, float learning_rate, float momentum, float decay)
{
    int size = layer.size*layer.size*layer.c*layer.n;
src/detection.c
@@ -3,11 +3,11 @@
#include "parser.h"
char *class_names[] = {"aeroplane", "bicycle", "bird", "boat", "bottle", "bus", "car", "cat", "chair", "cow", "diningtable", "dog", "horse", "motorbike", "person", "pottedplant", "sheep", "sofa", "train", "tvmonitor"};
char *class_names[] = {"bg", "aeroplane", "bicycle", "bird", "boat", "bottle", "bus", "car", "cat", "chair", "cow", "diningtable", "dog", "horse", "motorbike", "person", "pottedplant", "sheep", "sofa", "train", "tvmonitor"};
#define AMNT 3
void draw_detection(image im, float *box, int side)
{
    int classes = 20;
    int classes = 21;
    int elems = 4+classes;
    int j;
    int r, c;
@@ -50,6 +50,7 @@
    if(weightfile){
        load_weights(&net, weightfile);
    }
    net.seen = 0;
    printf("Learning Rate: %g, Momentum: %g, Decay: %g\n", net.learning_rate, net.momentum, net.decay);
    int imgs = 128;
    srand(time(0));
@@ -62,7 +63,7 @@
    int im_dim = 512;
    int jitter = 64;
    int classes = 20;
    int background = 1;
    int background = 0;
    pthread_t load_thread = load_data_detection_thread(imgs, paths, plist->size, classes, im_dim, im_dim, 7, 7, jitter, background, &buffer);
    clock_t time;
    while(1){
@@ -73,8 +74,8 @@
        load_thread = load_data_detection_thread(imgs, paths, plist->size, classes, im_dim, im_dim, 7, 7, jitter, background, &buffer);
        /*
           image im = float_to_image(im_dim - jitter, im_dim-jitter, 3, train.X.vals[0]);
           draw_detection(im, train.y.vals[0], 7);
           image im = float_to_image(im_dim - jitter, im_dim-jitter, 3, train.X.vals[114]);
           draw_detection(im, train.y.vals[114], 7);
           show_image(im, "truth");
           cvWaitKey(0);
         */
@@ -108,7 +109,7 @@
    char **paths = (char **)list_to_array(plist);
    int im_size = 448;
    int classes = 20;
    int background = 1;
    int background = 0;
    int num_output = 7*7*(4+classes+background);
    int m = plist->size;
src/dropout_layer_kernels.cu
@@ -11,7 +11,7 @@
    if(id < size) input[id] = (rand[id] < prob) ? 0 : input[id]*scale;
}
extern "C" void forward_dropout_layer_gpu(dropout_layer layer, network_state state)
void forward_dropout_layer_gpu(dropout_layer layer, network_state state)
{
    if (!state.train) return;
    int size = layer.inputs*layer.batch;
@@ -21,7 +21,7 @@
    check_error(cudaPeekAtLastError());
}
extern "C" void backward_dropout_layer_gpu(dropout_layer layer, network_state state)
void backward_dropout_layer_gpu(dropout_layer layer, network_state state)
{
    if(!state.delta) return;
    int size = layer.inputs*layer.batch;
src/network.c
@@ -194,24 +194,6 @@
    return get_network_delta_layer(net, net.n-1);
}
float calculate_error_network(network net, float *truth)
{
    float sum = 0;
    float *delta = get_network_delta(net);
    float *out = get_network_output(net);
    int i;
    for(i = 0; i < get_network_output_size(net)*net.batch; ++i){
        //if(i %get_network_output_size(net) == 0) printf("\n");
        //printf("%5.2f %5.2f, ", out[i], truth[i]);
        //if(i == get_network_output_size(net)) printf("\n");
        delta[i] = truth[i] - out[i];
        //printf("%.10f, ", out[i]);
        sum += delta[i]*delta[i];
    }
    //printf("\n");
    return sum;
}
int get_predicted_class_network(network net)
{
    float *out = get_network_output(net);
src/network_kernels.cu
@@ -20,8 +20,8 @@
#include "dropout_layer.h"
}
extern "C" float * get_network_output_gpu_layer(network net, int i);
extern "C" float * get_network_delta_gpu_layer(network net, int i);
float * get_network_output_gpu_layer(network net, int i);
float * get_network_delta_gpu_layer(network net, int i);
float *get_network_output_gpu(network net);
void forward_network_gpu(network net, network_state state)
@@ -196,8 +196,8 @@
    state.train = 1;
    forward_network_gpu(net, state);
    backward_network_gpu(net, state);
    if ((net.seen / net.batch) % net.subdivisions == 0) update_network_gpu(net);
    float error = get_network_cost(net);
    if ((net.seen / net.batch) % net.subdivisions == 0) update_network_gpu(net);
    return error;
}