diff --git a/src/convolutional_kernels.cu b/src/convolutional_kernels.cu index 9f0a2f85..a9a6837b 100644 --- a/src/convolutional_kernels.cu +++ b/src/convolutional_kernels.cu @@ -17,7 +17,7 @@ __global__ void bias_output_kernel(float *output, float *biases, int n, int size 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 @@ __global__ void backward_bias_kernel(float *bias_updates, float *delta, int batc } } -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<<>>(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 @@ extern "C" void forward_convolutional_layer_gpu(convolutional_layer layer, netwo 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 backward_convolutional_layer_gpu(convolutional_layer layer, netw } } -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 @@ extern "C" void pull_convolutional_layer(convolutional_layer layer) 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 @@ extern "C" void push_convolutional_layer(convolutional_layer layer) 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; diff --git a/src/detection.c b/src/detection.c index 15694c51..522a3219 100644 --- a/src/detection.c +++ b/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 @@ void train_detection(char *cfgfile, char *weightfile) 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 @@ void train_detection(char *cfgfile, char *weightfile) 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){ @@ -72,12 +73,12 @@ void train_detection(char *cfgfile, char *weightfile) train = buffer; 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); - */ +*/ printf("Loaded: %lf seconds\n", sec(clock()-time)); time=clock(); @@ -108,7 +109,7 @@ void validate_detection(char *cfgfile, char *weightfile) 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; diff --git a/src/dropout_layer_kernels.cu b/src/dropout_layer_kernels.cu index 4561d89d..2638ac50 100644 --- a/src/dropout_layer_kernels.cu +++ b/src/dropout_layer_kernels.cu @@ -11,7 +11,7 @@ __global__ void yoloswag420blazeit360noscope(float *input, int size, float *rand 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 @@ extern "C" void forward_dropout_layer_gpu(dropout_layer layer, network_state sta 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; diff --git a/src/network.c b/src/network.c index 61200d39..75c9454e 100644 --- a/src/network.c +++ b/src/network.c @@ -194,24 +194,6 @@ float *get_network_delta(network net) 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); diff --git a/src/network_kernels.cu b/src/network_kernels.cu index 4fc361d2..019f40d6 100644 --- a/src/network_kernels.cu +++ b/src/network_kernels.cu @@ -20,8 +20,8 @@ extern "C" { #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 @@ float train_network_datum_gpu(network net, float *x, float *y) 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; }