From 0f645836f193e75c4c3b718369e6fab15b5d19c5 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Wed, 11 Feb 2015 03:41:03 +0000
Subject: [PATCH] Detection is back, baby\!
---
src/network_kernels.cu | 38 +++++++++++++++++++++++++++++++++++++-
1 files changed, 37 insertions(+), 1 deletions(-)
diff --git a/src/network_kernels.cu b/src/network_kernels.cu
index a009174..1f3f2e0 100644
--- a/src/network_kernels.cu
+++ b/src/network_kernels.cu
@@ -10,6 +10,7 @@
#include "crop_layer.h"
#include "connected_layer.h"
#include "convolutional_layer.h"
+#include "deconvolutional_layer.h"
#include "maxpool_layer.h"
#include "cost_layer.h"
#include "normalization_layer.h"
@@ -31,6 +32,11 @@
forward_convolutional_layer_gpu(layer, input);
input = layer.output_gpu;
}
+ else if(net.types[i] == DECONVOLUTIONAL){
+ deconvolutional_layer layer = *(deconvolutional_layer *)net.layers[i];
+ forward_deconvolutional_layer_gpu(layer, input);
+ input = layer.output_gpu;
+ }
else if(net.types[i] == COST){
cost_layer layer = *(cost_layer *)net.layers[i];
forward_cost_layer_gpu(layer, input, truth);
@@ -58,9 +64,10 @@
}
else if(net.types[i] == CROP){
crop_layer layer = *(crop_layer *)net.layers[i];
- forward_crop_layer_gpu(layer, input);
+ forward_crop_layer_gpu(layer, train, input);
input = layer.output_gpu;
}
+ //cudaDeviceSynchronize();
//printf("Forward %d %s %f\n", i, get_layer_string(net.types[i]), sec(clock() - time));
}
}
@@ -83,6 +90,10 @@
convolutional_layer layer = *(convolutional_layer *)net.layers[i];
backward_convolutional_layer_gpu(layer, prev_input, prev_delta);
}
+ else if(net.types[i] == DECONVOLUTIONAL){
+ deconvolutional_layer layer = *(deconvolutional_layer *)net.layers[i];
+ backward_deconvolutional_layer_gpu(layer, prev_input, prev_delta);
+ }
else if(net.types[i] == COST){
cost_layer layer = *(cost_layer *)net.layers[i];
backward_cost_layer_gpu(layer, prev_input, prev_delta);
@@ -115,6 +126,10 @@
convolutional_layer layer = *(convolutional_layer *)net.layers[i];
update_convolutional_layer_gpu(layer);
}
+ else if(net.types[i] == DECONVOLUTIONAL){
+ deconvolutional_layer layer = *(deconvolutional_layer *)net.layers[i];
+ update_deconvolutional_layer_gpu(layer);
+ }
else if(net.types[i] == CONNECTED){
connected_layer layer = *(connected_layer *)net.layers[i];
update_connected_layer_gpu(layer);
@@ -128,6 +143,10 @@
convolutional_layer layer = *(convolutional_layer *)net.layers[i];
return layer.output_gpu;
}
+ else if(net.types[i] == DECONVOLUTIONAL){
+ deconvolutional_layer layer = *(deconvolutional_layer *)net.layers[i];
+ return layer.output_gpu;
+ }
else if(net.types[i] == CONNECTED){
connected_layer layer = *(connected_layer *)net.layers[i];
return layer.output_gpu;
@@ -156,6 +175,10 @@
convolutional_layer layer = *(convolutional_layer *)net.layers[i];
return layer.delta_gpu;
}
+ else if(net.types[i] == DECONVOLUTIONAL){
+ deconvolutional_layer layer = *(deconvolutional_layer *)net.layers[i];
+ return layer.delta_gpu;
+ }
else if(net.types[i] == CONNECTED){
connected_layer layer = *(connected_layer *)net.layers[i];
return layer.delta_gpu;
@@ -176,6 +199,7 @@
float train_network_datum_gpu(network net, float *x, float *y)
{
+ //clock_t time = clock();
int x_size = get_network_input_size(net)*net.batch;
int y_size = get_network_output_size(net)*net.batch;
if(!*net.input_gpu){
@@ -185,10 +209,18 @@
cuda_push_array(*net.input_gpu, x, x_size);
cuda_push_array(*net.truth_gpu, y, y_size);
}
+ //printf("trans %f\n", sec(clock() - time));
+ //time = clock();
forward_network_gpu(net, *net.input_gpu, *net.truth_gpu, 1);
+ //printf("forw %f\n", sec(clock() - time));
+ //time = clock();
backward_network_gpu(net, *net.input_gpu);
+ //printf("back %f\n", sec(clock() - time));
+ //time = clock();
update_network_gpu(net);
float error = get_network_cost(net);
+ //printf("updt %f\n", sec(clock() - time));
+ //time = clock();
return error;
}
@@ -198,6 +230,10 @@
convolutional_layer layer = *(convolutional_layer *)net.layers[i];
return layer.output;
}
+ else if(net.types[i] == DECONVOLUTIONAL){
+ deconvolutional_layer layer = *(deconvolutional_layer *)net.layers[i];
+ return layer.output;
+ }
else if(net.types[i] == CONNECTED){
connected_layer layer = *(connected_layer *)net.layers[i];
cuda_pull_array(layer.output_gpu, layer.output, layer.outputs*layer.batch);
--
Gitblit v1.10.0