rolling avg demo

cfg/darknet.cfg

@@ -7,11 +7,10 @@ channels=3
 momentum=0.9
 decay=0.0005
 
-learning_rate=0.01
+learning_rate=0.1
-policy=sigmoid
+policy=poly
-gamma=.00002
+power=4
-step=400000
+max_batches=500000
-max_batches=800000
 
 [crop]
 crop_height=224
@@ -22,6 +21,7 @@ saturation=1
 exposure=1
 
 [convolutional]
+batch_normalize=1
 filters=16
 size=3
 stride=1
@@ -33,6 +33,7 @@ size=2
 stride=2
 
 [convolutional]
+batch_normalize=1
 filters=32
 size=3
 stride=1
@@ -44,6 +45,7 @@ size=2
 stride=2
 
 [convolutional]
+batch_normalize=1
 filters=64
 size=3
 stride=1
@@ -55,6 +57,7 @@ size=2
 stride=2
 
 [convolutional]
+batch_normalize=1
 filters=128
 size=3
 stride=1
@@ -66,6 +69,7 @@ size=2
 stride=2
 
 [convolutional]
+batch_normalize=1
 filters=256
 size=3
 stride=1
@@ -77,6 +81,7 @@ size=2
 stride=2
 
 [convolutional]
+batch_normalize=1
 filters=512
 size=3
 stride=1
@@ -88,6 +93,7 @@ size=2
 stride=2
 
 [convolutional]
+batch_normalize=1
 filters=1024
 size=3
 stride=1

src/coco.c

@@ -385,11 +385,15 @@ void test_coco(char *cfgfile, char *weightfile, char *filename, float thresh)
     }
 }
 
-#ifdef OPENCV
+void demo_coco(char *cfgfile, char *weightfile, float thresh, int cam_index, char *filename);
-#ifdef GPU
+static void demo(char *cfgfile, char *weightfile, float thresh, int cam_index, char* filename)
-void demo_coco(char *cfgfile, char *weightfile, float thresh, int cam_index);
+{
-#endif
+    #if defined(OPENCV) && defined(GPU)
-#endif
+        demo_coco(cfgfile, weightfile, thresh, cam_index, filename);
+    #else
+        fprintf(stderr, "Need to compile with GPU and OpenCV for demo.\n");
+    #endif
+}
 
 void run_coco(int argc, char **argv)
 {
@@ -401,6 +405,7 @@ void run_coco(int argc, char **argv)
     }
     float thresh = find_float_arg(argc, argv, "-thresh", .2);
     int cam_index = find_int_arg(argc, argv, "-c", 0);
+    char *file = find_char_arg(argc, argv, "-file", 0);
 
     if(argc < 4){
         fprintf(stderr, "usage: %s %s [train/test/valid] [cfg] [weights (optional)]\n", argv[0], argv[1]);
@@ -414,9 +419,5 @@ void run_coco(int argc, char **argv)
     else if(0==strcmp(argv[2], "train")) train_coco(cfg, weights);
     else if(0==strcmp(argv[2], "valid")) validate_coco(cfg, weights);
     else if(0==strcmp(argv[2], "recall")) validate_coco_recall(cfg, weights);
-#ifdef OPENCV
+    else if(0==strcmp(argv[2], "demo")) demo(cfg, weights, thresh, cam_index, file);
-#ifdef GPU
-    else if(0==strcmp(argv[2], "demo")) demo_coco(cfg, weights, thresh, cam_index);
-#endif
-#endif
 }

src/coco_kernels.cu

@@ -34,6 +34,12 @@ static cv::VideoCapture cap;
 static float fps = 0;
 static float demo_thresh = 0;
 
+static const int frames = 3;
+static float *predictions[frames];
+static int demo_index = 0;
+static image images[frames];
+static float *avg;
+
 void *fetch_in_thread_coco(void *ptr)
 {
     cv::Mat frame_m;
@@ -51,19 +57,28 @@ void *detect_in_thread_coco(void *ptr)
 
     detection_layer l = net.layers[net.n-1];
     float *X = det_s.data;
-    float *predictions = network_predict(net, X);
+    float *prediction = network_predict(net, X);
+
+    memcpy(predictions[demo_index], prediction, l.outputs*sizeof(float));
+    mean_arrays(predictions, frames, l.outputs, avg);
+
     free_image(det_s);
-    convert_coco_detections(predictions, l.classes, l.n, l.sqrt, l.side, 1, 1, demo_thresh, probs, boxes, 0);
+    convert_coco_detections(avg, l.classes, l.n, l.sqrt, l.side, 1, 1, demo_thresh, probs, boxes, 0);
     if (nms > 0) do_nms(boxes, probs, l.side*l.side*l.n, l.classes, nms);
     printf("\033[2J");
     printf("\033[1;1H");
     printf("\nFPS:%.0f\n",fps);
     printf("Objects:\n\n");
+
+    images[demo_index] = det;
+    det = images[(demo_index + frames/2 + 1)%frames];
+    demo_index = (demo_index + 1)%frames;
+
     draw_detections(det, l.side*l.side*l.n, demo_thresh, boxes, probs, coco_classes, coco_labels, 80);
     return 0;
 }
 
-extern "C" void demo_coco(char *cfgfile, char *weightfile, float thresh, int cam_index)
+extern "C" void demo_coco(char *cfgfile, char *weightfile, float thresh, int cam_index, const char *filename)
 {
     demo_thresh = thresh;
     printf("YOLO demo\n");
@@ -75,13 +90,21 @@ extern "C" void demo_coco(char *cfgfile, char *weightfile, float thresh, int cam
 
     srand(2222222);
 
-    cv::VideoCapture cam(cam_index);
+    if(filename){
-    cap = cam;
+        cap.open(filename);
+    }else{
+        cap.open(cam_index);
+    }
+
     if(!cap.isOpened()) error("Couldn't connect to webcam.\n");
 
     detection_layer l = net.layers[net.n-1];
     int j;
 
+    avg = (float *) calloc(l.outputs, sizeof(float));
+    for(j = 0; j < frames; ++j) predictions[j] = (float *) calloc(l.outputs, sizeof(float));
+    for(j = 0; j < frames; ++j) images[j] = make_image(1,1,3);
+
     boxes = (box *)calloc(l.side*l.side*l.n, sizeof(box));
     probs = (float **)calloc(l.side*l.side*l.n, sizeof(float *));
     for(j = 0; j < l.side*l.side*l.n; ++j) probs[j] = (float *)calloc(l.classes, sizeof(float *));

src/local_kernels.cu

@@ -1,230 +0,0 @@
-#include "cuda_runtime.h"
-#include "curand.h"
-#include "cublas_v2.h"
-
-extern "C" {
-#include "local_layer.h"
-#include "gemm.h"
-#include "blas.h"
-#include "im2col.h"
-#include "col2im.h"
-#include "utils.h"
-#include "cuda.h"
-}
-
-__global__ void scale_bias_kernel(float *output, float *biases, int n, int size)
-{
-    int offset = blockIdx.x * blockDim.x + threadIdx.x;
-    int filter = blockIdx.y;
-    int batch = blockIdx.z;
-
-    if(offset < size) output[(batch*n+filter)*size + offset] *= biases[filter];
-}
-
-void scale_bias_gpu(float *output, float *biases, int batch, int n, int size)
-{
-    dim3 dimGrid((size-1)/BLOCK + 1, n, batch);
-    dim3 dimBlock(BLOCK, 1, 1);
-
-    scale_bias_kernel<<<dimGrid, dimBlock>>>(output, biases, n, size);
-    check_error(cudaPeekAtLastError());
-}
-
-__global__ void backward_scale_kernel(float *x_norm, float *delta, int batch, int n, int size, float *scale_updates)
-{
-    __shared__ float part[BLOCK];
-    int i,b;
-    int filter = blockIdx.x;
-    int p = threadIdx.x;
-    float sum = 0;
-    for(b = 0; b < batch; ++b){
-        for(i = 0; i < size; i += BLOCK){
-            int index = p + i + size*(filter + n*b);
-            sum += (p+i < size) ? delta[index]*x_norm[index] : 0;
-        }
-    }
-    part[p] = sum;
-    __syncthreads();
-    if (p == 0) {
-        for(i = 0; i < BLOCK; ++i) scale_updates[filter] += part[i];
-    }
-}
-
-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);
-    check_error(cudaPeekAtLastError());
-}
-
-__global__ void add_bias_kernel(float *output, float *biases, int n, int size)
-{
-    int offset = blockIdx.x * blockDim.x + threadIdx.x;
-    int filter = blockIdx.y;
-    int batch = blockIdx.z;
-
-    if(offset < size) output[(batch*n+filter)*size + offset] += biases[filter];
-}
-
-void add_bias_gpu(float *output, float *biases, int batch, int n, int size)
-{
-    dim3 dimGrid((size-1)/BLOCK + 1, n, batch);
-    dim3 dimBlock(BLOCK, 1, 1);
-
-    add_bias_kernel<<<dimGrid, dimBlock>>>(output, biases, n, size);
-    check_error(cudaPeekAtLastError());
-}
-
-__global__ void backward_bias_kernel(float *bias_updates, float *delta, int batch, int n, int size)
-{
-    __shared__ float part[BLOCK];
-    int i,b;
-    int filter = blockIdx.x;
-    int p = threadIdx.x;
-    float sum = 0;
-    for(b = 0; b < batch; ++b){
-        for(i = 0; i < size; i += BLOCK){
-            int index = p + i + size*(filter + n*b);
-            sum += (p+i < size) ? delta[index] : 0;
-        }
-    }
-    part[p] = sum;
-    __syncthreads();
-    if (p == 0) {
-        for(i = 0; i < BLOCK; ++i) bias_updates[filter] += part[i];
-    }
-}
-
-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);
-    check_error(cudaPeekAtLastError());
-}
-
-void forward_local_layer_gpu(local_layer l, network_state state)
-{
-    int i;
-    int m = l.n;
-    int k = l.size*l.size*l.c;
-    int n = local_out_height(l)*
-        local_out_width(l);
-
-    fill_ongpu(l.outputs*l.batch, 0, l.output_gpu, 1);
-    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;
-        float * b = l.col_image_gpu;
-        float * c = l.output_gpu;
-        gemm_ongpu(0,0,m,n,k,1.,a,k,b,n,1.,c+i*m*n,n);
-    }
-
-    if(l.batch_normalize){
-        if(state.train){
-            fast_mean_gpu(l.output_gpu, l.batch, l.n, l.out_h*l.out_w, l.spatial_mean_gpu, l.mean_gpu);   
-            fast_variance_gpu(l.output_gpu, l.mean_gpu, l.batch, l.n, l.out_h*l.out_w, l.spatial_variance_gpu, l.variance_gpu);
-
-            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);
-            axpy_ongpu(l.n, .05, l.variance_gpu, 1, l.rolling_variance_gpu, 1);
-
-            // cuda_pull_array(l.variance_gpu, l.mean, l.n);
-            // printf("%f\n", l.mean[0]);
-
-            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.n, l.out_h*l.out_w);
-            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.n, l.out_h*l.out_w);
-        }
-
-        scale_bias_gpu(l.output_gpu, l.scales_gpu, l.batch, l.n, l.out_h*l.out_w);
-    }
-    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);
-}
-
-void backward_local_layer_gpu(local_layer l, network_state state)
-{
-    int i;
-    int m = l.n;
-    int n = l.size*l.size*l.c;
-    int k = local_out_height(l)*
-        local_out_width(l);
-
-    gradient_array_ongpu(l.output_gpu, m*k*l.batch, l.activation, l.delta_gpu);
-
-    backward_bias_gpu(l.bias_updates_gpu, l.delta_gpu, l.batch, l.n, k);
-
-    if(l.batch_normalize){
-        backward_scale_gpu(l.x_norm_gpu, l.delta_gpu, l.batch, l.n, l.out_w*l.out_h, l.scale_updates_gpu);
-
-        scale_bias_gpu(l.delta_gpu, l.scales_gpu, l.batch, l.n, l.out_h*l.out_w);
-
-        fast_mean_delta_gpu(l.delta_gpu, l.variance_gpu, l.batch, l.n, l.out_w*l.out_h, l.spatial_mean_delta_gpu, l.mean_delta_gpu);
-        fast_variance_delta_gpu(l.x_gpu, l.delta_gpu, l.mean_gpu, l.variance_gpu, l.batch, l.n, l.out_w*l.out_h, l.spatial_variance_delta_gpu, 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.n, l.out_w*l.out_h, l.delta_gpu);
-    }
-
-    for(i = 0; i < l.batch; ++i){
-        float * a = l.delta_gpu;
-        float * b = l.col_image_gpu;
-        float * c = l.filter_updates_gpu;
-
-        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);
-        gemm_ongpu(0,1,m,n,k,1,a + i*m*k,k,b,k,1,c,n);
-
-        if(state.delta){
-            float * a = l.filters_gpu;
-            float * b = l.delta_gpu;
-            float * c = l.col_image_gpu;
-
-            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);
-        }
-    }
-}
-
-void pull_local_layer(local_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);
-    cuda_pull_array(layer.filter_updates_gpu, layer.filter_updates, layer.c*layer.n*layer.size*layer.size);
-    cuda_pull_array(layer.bias_updates_gpu, layer.bias_updates, layer.n);
-    if (layer.batch_normalize){
-        cuda_pull_array(layer.scales_gpu, layer.scales, layer.n);
-        cuda_pull_array(layer.rolling_mean_gpu, layer.rolling_mean, layer.n);
-        cuda_pull_array(layer.rolling_variance_gpu, layer.rolling_variance, layer.n);
-    }
-}
-
-void push_local_layer(local_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);
-    cuda_push_array(layer.filter_updates_gpu, layer.filter_updates, layer.c*layer.n*layer.size*layer.size);
-    cuda_push_array(layer.bias_updates_gpu, layer.bias_updates, layer.n);
-    if (layer.batch_normalize){
-        cuda_push_array(layer.scales_gpu, layer.scales, layer.n);
-        cuda_push_array(layer.rolling_mean_gpu, layer.rolling_mean, layer.n);
-        cuda_push_array(layer.rolling_variance_gpu, layer.rolling_variance, layer.n);
-    }
-}
-
-void update_local_layer_gpu(local_layer layer, int batch, float learning_rate, float momentum, float decay)
-{
-    int size = layer.size*layer.size*layer.c*layer.n;
-
-    axpy_ongpu(layer.n, learning_rate/batch, layer.bias_updates_gpu, 1, layer.biases_gpu, 1);
-    scal_ongpu(layer.n, momentum, layer.bias_updates_gpu, 1);
-
-    axpy_ongpu(layer.n, learning_rate/batch, layer.scale_updates_gpu, 1, layer.scales_gpu, 1);
-    scal_ongpu(layer.n, momentum, layer.scale_updates_gpu, 1);
-
-    axpy_ongpu(size, -decay*batch, layer.filters_gpu, 1, layer.filter_updates_gpu, 1);
-    axpy_ongpu(size, learning_rate/batch, layer.filter_updates_gpu, 1, layer.filters_gpu, 1);
-    scal_ongpu(size, momentum, layer.filter_updates_gpu, 1);
-}
-
-

src/utils.c

@@ -359,6 +359,21 @@ float mean_array(float *a, int n)
     return sum_array(a,n)/n;
 }
 
+void mean_arrays(float **a, int n, int els, float *avg)
+{
+    int i;
+    int j;
+    memset(avg, 0, els*sizeof(float));
+    for(j = 0; j < n; ++j){
+        for(i = 0; i < els; ++i){
+            avg[i] += a[j][i];
+        }
+    }
+    for(i = 0; i < els; ++i){
+        avg[i] /= n;
+    }
+}
+
 float variance_array(float *a, int n)
 {
     int i;

src/utils.h

@@ -37,6 +37,7 @@ float rand_normal();
 float rand_uniform();
 float sum_array(float *a, int n);
 float mean_array(float *a, int n);
+void mean_arrays(float **a, int n, int els, float *avg);
 float variance_array(float *a, int n);
 float mag_array(float *a, int n);
 float **one_hot_encode(float *a, int n, int k);