| | |
| | | #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;} |
| | |
| | | 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) |
| | |
| | | return plse_activate_kernel(x); |
| | | case STAIR: |
| | | return stair_activate_kernel(x); |
| | | case HARDTAN: |
| | | return hardtan_activate_kernel(x); |
| | | } |
| | | return 0; |
| | | } |
| | |
| | | return plse_gradient_kernel(x); |
| | | case STAIR: |
| | | return stair_gradient_kernel(x); |
| | | case HARDTAN: |
| | | return hardtan_gradient_kernel(x); |
| | | } |
| | | return 0; |
| | | } |
| | |
| | | return "leaky"; |
| | | case STAIR: |
| | | return "stair"; |
| | | case HARDTAN: |
| | | return "hardtan"; |
| | | default: |
| | | break; |
| | | } |
| | |
| | | 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; |
| | |
| | | return plse_activate(x); |
| | | case STAIR: |
| | | return stair_activate(x); |
| | | case HARDTAN: |
| | | return hardtan_activate(x); |
| | | } |
| | | return 0; |
| | | } |
| | |
| | | return plse_gradient(x); |
| | | case STAIR: |
| | | return stair_gradient(x); |
| | | case HARDTAN: |
| | | return hardtan_gradient(x); |
| | | } |
| | | return 0; |
| | | } |
| | |
| | | #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); |
| | |
| | | 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;} |
| | |
| | | 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) |
| | |
| | | int *indexes = calloc(top, sizeof(int)); |
| | | char buff[256]; |
| | | char *input = buff; |
| | | int size = net.w; |
| | | while(1){ |
| | | if(filename){ |
| | | strncpy(input, filename, 256); |
| | |
| | | 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); |
| | |
| | | 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; |
| | | } |
| | |
| | | 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 |
| | |
| | | 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 |
| | |
| | | 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 |
| | |
| | | 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; |
| | |
| | | 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; |
| | | } |