This commit is contained in:
Joseph Redmon 2016-06-09 17:20:31 -07:00
parent 32d2c96997
commit 729ce43e6e
7 changed files with 52 additions and 6 deletions

View File

@ -7,6 +7,13 @@ extern "C" {
#include "cuda.h"
}
__device__ float hardtan_activate_kernel(float x)
{
if (x < -1) return -1;
if (x > 1) return 1;
return x;
}
__device__ float linear_activate_kernel(float x){return x;}
__device__ float logistic_activate_kernel(float x){return 1./(1. + exp(-x));}
__device__ float loggy_activate_kernel(float x){return 2./(1. + exp(-x)) - 1;}
@ -29,6 +36,12 @@ __device__ float stair_activate_kernel(float x)
else return (x - n) + floor(x/2.);
}
__device__ float hardtan_gradient_kernel(float x)
{
if (x > -1 && x < 1) return 1;
return 0;
}
__device__ float linear_gradient_kernel(float x){return 1;}
__device__ float logistic_gradient_kernel(float x){return (1-x)*x;}
__device__ float loggy_gradient_kernel(float x)
@ -74,6 +87,8 @@ __device__ float activate_kernel(float x, ACTIVATION a)
return plse_activate_kernel(x);
case STAIR:
return stair_activate_kernel(x);
case HARDTAN:
return hardtan_activate_kernel(x);
}
return 0;
}
@ -103,6 +118,8 @@ __device__ float gradient_kernel(float x, ACTIVATION a)
return plse_gradient_kernel(x);
case STAIR:
return stair_gradient_kernel(x);
case HARDTAN:
return hardtan_gradient_kernel(x);
}
return 0;
}

View File

@ -30,6 +30,8 @@ char *get_activation_string(ACTIVATION a)
return "leaky";
case STAIR:
return "stair";
case HARDTAN:
return "hardtan";
default:
break;
}
@ -44,6 +46,7 @@ ACTIVATION get_activation(char *s)
if (strcmp(s, "elu")==0) return ELU;
if (strcmp(s, "relie")==0) return RELIE;
if (strcmp(s, "plse")==0) return PLSE;
if (strcmp(s, "hardtan")==0) return HARDTAN;
if (strcmp(s, "linear")==0) return LINEAR;
if (strcmp(s, "ramp")==0) return RAMP;
if (strcmp(s, "leaky")==0) return LEAKY;
@ -78,6 +81,8 @@ float activate(float x, ACTIVATION a)
return plse_activate(x);
case STAIR:
return stair_activate(x);
case HARDTAN:
return hardtan_activate(x);
}
return 0;
}
@ -115,6 +120,8 @@ float gradient(float x, ACTIVATION a)
return plse_gradient(x);
case STAIR:
return stair_gradient(x);
case HARDTAN:
return hardtan_gradient(x);
}
return 0;
}

View File

@ -4,7 +4,7 @@
#include "math.h"
typedef enum{
LOGISTIC, RELU, RELIE, LINEAR, RAMP, TANH, PLSE, LEAKY, ELU, LOGGY, STAIR
LOGISTIC, RELU, RELIE, LINEAR, RAMP, TANH, PLSE, LEAKY, ELU, LOGGY, STAIR, HARDTAN
}ACTIVATION;
ACTIVATION get_activation(char *s);
@ -25,6 +25,12 @@ static inline float stair_activate(float x)
if (n%2 == 0) return floor(x/2.);
else return (x - n) + floor(x/2.);
}
static inline float hardtan_activate(float x)
{
if (x < -1) return -1;
if (x > 1) return 1;
return x;
}
static inline float linear_activate(float x){return x;}
static inline float logistic_activate(float x){return 1./(1. + exp(-x));}
static inline float loggy_activate(float x){return 2./(1. + exp(-x)) - 1;}
@ -41,6 +47,11 @@ static inline float plse_activate(float x)
return .125*x + .5;
}
static inline float hardtan_gradient(float x)
{
if (x > -1 && x < 1) return 1;
return 0;
}
static inline float linear_gradient(float x){return 1;}
static inline float logistic_gradient(float x){return (1-x)*x;}
static inline float loggy_gradient(float x)

View File

@ -477,6 +477,7 @@ void predict_classifier(char *datacfg, char *cfgfile, char *weightfile, char *fi
int *indexes = calloc(top, sizeof(int));
char buff[256];
char *input = buff;
int size = net.w;
while(1){
if(filename){
strncpy(input, filename, 256);
@ -487,8 +488,12 @@ void predict_classifier(char *datacfg, char *cfgfile, char *weightfile, char *fi
if(!input) return;
strtok(input, "\n");
}
image im = load_image_color(input, net.w, net.h);
float *X = im.data;
image im = load_image_color(input, 0, 0);
image r = resize_min(im, size);
resize_network(&net, r.w, r.h);
printf("%d %d\n", r.w, r.h);
float *X = r.data;
time=clock();
float *predictions = network_predict(net, X);
top_predictions(net, top, indexes);
@ -497,6 +502,7 @@ void predict_classifier(char *datacfg, char *cfgfile, char *weightfile, char *fi
int index = indexes[i];
printf("%s: %f\n", names[index], predictions[index]);
}
if(r.data != im.data) free_image(r);
free_image(im);
if (filename) break;
}

View File

@ -142,6 +142,7 @@ void backward_convolutional_layer_gpu(convolutional_layer l, network_state state
if(l.batch_normalize){
backward_batchnorm_layer_gpu(l, state);
}
float *original_input = state.input;
if(l.xnor) state.input = l.binary_input_gpu;
#ifdef CUDNN
@ -176,6 +177,7 @@ void backward_convolutional_layer_gpu(convolutional_layer l, network_state state
l.dsrcTensorDesc,
state.delta);
if(l.binary || l.xnor) swap_binary(&l);
if(l.xnor) gradient_array_ongpu(original_input, l.batch*l.c*l.h*l.w, HARDTAN, state.delta);
}
#else
@ -197,7 +199,10 @@ void backward_convolutional_layer_gpu(convolutional_layer l, network_state state
gemm_ongpu(1,0,n,k,m,1,a,n,b + i*k*m,k,0,c,k);
col2im_ongpu(state.workspace, l.c, l.h, l.w, l.size, l.stride, l.pad, state.delta + i*l.c*l.h*l.w);
if(l.binary || l.xnor) swap_binary(&l);
if(l.binary || l.xnor) {
swap_binary(&l);
}
if(l.xnor) gradient_array_ongpu(original_input + i*l.c*l.h*l.w, l.c*l.h*l.w, HARDTAN, state.delta + i*l.c*l.h*l.w);
}
}
#endif

View File

@ -434,7 +434,7 @@ int resize_network(network *net, int w, int h)
net->workspace = cuda_make_array(0, (workspace_size-1)/sizeof(float)+1);
#else
free(net->workspace);
net->workspace = calloc(1, (workspace_size-1)/sizeof(float)+1);
net->workspace = calloc(1, workspace_size);
#endif
//fprintf(stderr, " Done!\n");
return 0;

View File

@ -125,7 +125,7 @@ void test_tag(char *cfgfile, char *weightfile, char *filename)
int index = indexes[i];
printf("%.1f%%: %s\n", predictions[index]*100, names[index]);
}
free_image(r);
if(r.data != im.data) free_image(r);
free_image(im);
if (filename) break;
}