From 73f7aacf35ec9b1d0f9de9ddf38af0889f213e99 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Tue, 20 Sep 2016 18:34:49 +0000
Subject: [PATCH] better multigpu
---
src/network.c | 105 +++++++++++++++++++++++++++++++++++++++++++++++-----
1 files changed, 95 insertions(+), 10 deletions(-)
diff --git a/src/network.c b/src/network.c
index 8dee8cc..72c8943 100644
--- a/src/network.c
+++ b/src/network.c
@@ -1,5 +1,6 @@
#include <stdio.h>
#include <time.h>
+#include <assert.h>
#include "network.h"
#include "image.h"
#include "data.h"
@@ -8,12 +9,19 @@
#include "crop_layer.h"
#include "connected_layer.h"
+#include "gru_layer.h"
+#include "rnn_layer.h"
+#include "crnn_layer.h"
#include "local_layer.h"
#include "convolutional_layer.h"
+#include "activation_layer.h"
#include "deconvolutional_layer.h"
#include "detection_layer.h"
+#include "region_layer.h"
#include "normalization_layer.h"
+#include "batchnorm_layer.h"
#include "maxpool_layer.h"
+#include "reorg_layer.h"
#include "avgpool_layer.h"
#include "cost_layer.h"
#include "softmax_layer.h"
@@ -59,7 +67,10 @@
case EXP:
return net.learning_rate * pow(net.gamma, batch_num);
case POLY:
+ if (batch_num < net.burn_in) return net.learning_rate * pow((float)batch_num / net.burn_in, net.power);
return net.learning_rate * pow(1 - (float)batch_num / net.max_batches, net.power);
+ case RANDOM:
+ return net.learning_rate * pow(rand_uniform(0,1), net.power);
case SIG:
return net.learning_rate * (1./(1.+exp(net.gamma*(batch_num - net.step))));
default:
@@ -73,20 +84,32 @@
switch(a){
case CONVOLUTIONAL:
return "convolutional";
+ case ACTIVE:
+ return "activation";
case LOCAL:
return "local";
case DECONVOLUTIONAL:
return "deconvolutional";
case CONNECTED:
return "connected";
+ case RNN:
+ return "rnn";
+ case GRU:
+ return "gru";
+ case CRNN:
+ return "crnn";
case MAXPOOL:
return "maxpool";
+ case REORG:
+ return "reorg";
case AVGPOOL:
return "avgpool";
case SOFTMAX:
return "softmax";
case DETECTION:
return "detection";
+ case REGION:
+ return "region";
case DROPOUT:
return "dropout";
case CROP:
@@ -99,6 +122,8 @@
return "shortcut";
case NORMALIZATION:
return "normalization";
+ case BATCHNORM:
+ return "batchnorm";
default:
break;
}
@@ -120,6 +145,7 @@
void forward_network(network net, network_state state)
{
+ state.workspace = net.workspace;
int i;
for(i = 0; i < net.n; ++i){
state.index = i;
@@ -131,14 +157,26 @@
forward_convolutional_layer(l, state);
} else if(l.type == DECONVOLUTIONAL){
forward_deconvolutional_layer(l, state);
+ } else if(l.type == ACTIVE){
+ forward_activation_layer(l, state);
} else if(l.type == LOCAL){
forward_local_layer(l, state);
} else if(l.type == NORMALIZATION){
forward_normalization_layer(l, state);
+ } else if(l.type == BATCHNORM){
+ forward_batchnorm_layer(l, state);
} else if(l.type == DETECTION){
forward_detection_layer(l, state);
+ } else if(l.type == REGION){
+ forward_region_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 == GRU){
+ forward_gru_layer(l, state);
+ } else if(l.type == CRNN){
+ forward_crnn_layer(l, state);
} else if(l.type == CROP){
forward_crop_layer(l, state);
} else if(l.type == COST){
@@ -147,6 +185,8 @@
forward_softmax_layer(l, state);
} else if(l.type == MAXPOOL){
forward_maxpool_layer(l, state);
+ } else if(l.type == REORG){
+ forward_reorg_layer(l, state);
} else if(l.type == AVGPOOL){
forward_avgpool_layer(l, state);
} else if(l.type == DROPOUT){
@@ -173,6 +213,12 @@
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 == GRU){
+ update_gru_layer(l, update_batch, rate, net.momentum, net.decay);
+ } else if(l.type == CRNN){
+ update_crnn_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);
}
@@ -181,6 +227,9 @@
float *get_network_output(network net)
{
+ #ifdef GPU
+ if (gpu_index >= 0) return get_network_output_gpu(net);
+ #endif
int i;
for(i = net.n-1; i > 0; --i) if(net.layers[i].type != COST) break;
return net.layers[i].output;
@@ -192,11 +241,7 @@
float sum = 0;
int count = 0;
for(i = 0; i < net.n; ++i){
- if(net.layers[i].type == COST){
- sum += net.layers[i].output[0];
- ++count;
- }
- if(net.layers[i].type == DETECTION){
+ if(net.layers[i].cost){
sum += net.layers[i].cost[0];
++count;
}
@@ -216,6 +261,7 @@
int i;
float *original_input = state.input;
float *original_delta = state.delta;
+ state.workspace = net.workspace;
for(i = net.n-1; i >= 0; --i){
state.index = i;
if(i == 0){
@@ -231,20 +277,34 @@
backward_convolutional_layer(l, state);
} else if(l.type == DECONVOLUTIONAL){
backward_deconvolutional_layer(l, state);
+ } else if(l.type == ACTIVE){
+ backward_activation_layer(l, state);
} else if(l.type == NORMALIZATION){
backward_normalization_layer(l, state);
+ } else if(l.type == BATCHNORM){
+ backward_batchnorm_layer(l, state);
} else if(l.type == MAXPOOL){
if(i != 0) backward_maxpool_layer(l, state);
+ } else if(l.type == REORG){
+ backward_reorg_layer(l, state);
} else if(l.type == AVGPOOL){
backward_avgpool_layer(l, state);
} else if(l.type == DROPOUT){
backward_dropout_layer(l, state);
} else if(l.type == DETECTION){
backward_detection_layer(l, state);
+ } else if(l.type == REGION){
+ backward_region_layer(l, state);
} else if(l.type == SOFTMAX){
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 == GRU){
+ backward_gru_layer(l, state);
+ } else if(l.type == CRNN){
+ backward_crnn_layer(l, state);
} else if(l.type == LOCAL){
backward_local_layer(l, state);
} else if(l.type == COST){
@@ -259,11 +319,11 @@
float train_network_datum(network net, float *x, float *y)
{
- *net.seen += net.batch;
#ifdef GPU
if(gpu_index >= 0) return train_network_datum_gpu(net, x, y);
#endif
network_state state;
+ *net.seen += net.batch;
state.index = 0;
state.net = net;
state.input = x;
@@ -297,6 +357,7 @@
float train_network(network net, data d)
{
+ assert(d.X.rows % net.batch == 0);
int batch = net.batch;
int n = d.X.rows / batch;
float *X = calloc(batch*d.X.cols, sizeof(float));
@@ -314,6 +375,7 @@
return (float)sum/(n*batch);
}
+
float train_network_batch(network net, data d, int n)
{
int i,j;
@@ -344,6 +406,11 @@
int i;
for(i = 0; i < net->n; ++i){
net->layers[i].batch = b;
+ #ifdef CUDNN
+ if(net->layers[i].type == CONVOLUTIONAL){
+ cudnn_convolutional_setup(net->layers + i);
+ }
+ #endif
}
}
@@ -354,17 +421,21 @@
net->w = w;
net->h = h;
int inputs = 0;
- //fprintf(stderr, "Resizing to %d x %d...", w, h);
+ size_t workspace_size = 0;
+ //fprintf(stderr, "Resizing to %d x %d...\n", w, h);
//fflush(stderr);
for (i = 0; i < net->n; ++i){
layer l = net->layers[i];
if(l.type == CONVOLUTIONAL){
resize_convolutional_layer(&l, w, h);
+ }else if(l.type == CROP){
+ resize_crop_layer(&l, w, h);
}else if(l.type == MAXPOOL){
resize_maxpool_layer(&l, w, h);
+ }else if(l.type == REORG){
+ resize_reorg_layer(&l, w, h);
}else if(l.type == AVGPOOL){
resize_avgpool_layer(&l, w, h);
- break;
}else if(l.type == NORMALIZATION){
resize_normalization_layer(&l, w, h);
}else if(l.type == COST){
@@ -372,11 +443,25 @@
}else{
error("Cannot resize this type of layer");
}
+ if(l.workspace_size > workspace_size) workspace_size = l.workspace_size;
inputs = l.outputs;
net->layers[i] = l;
w = l.out_w;
h = l.out_h;
+ if(l.type == AVGPOOL) break;
}
+#ifdef GPU
+ if(gpu_index >= 0){
+ cuda_free(net->workspace);
+ net->workspace = cuda_make_array(0, (workspace_size-1)/sizeof(float)+1);
+ }else {
+ free(net->workspace);
+ net->workspace = calloc(1, workspace_size);
+ }
+#else
+ free(net->workspace);
+ net->workspace = calloc(1, workspace_size);
+#endif
//fprintf(stderr, " Done!\n");
return 0;
}
@@ -591,10 +676,10 @@
free_layer(net.layers[i]);
}
free(net.layers);
- #ifdef GPU
+#ifdef GPU
if(*net.input_gpu) cuda_free(*net.input_gpu);
if(*net.truth_gpu) cuda_free(*net.truth_gpu);
if(net.input_gpu) free(net.input_gpu);
if(net.truth_gpu) free(net.truth_gpu);
- #endif
+#endif
}
--
Gitblit v1.10.0