So there WAS this huge bug. Gone now

Joseph Redmon

2014-05-09 cd8d53df21f3ad2810add2a8cff766c745f55a17

 src/activations.c

@@ -2,6 +2,7 @@

#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

char *get_activation_string(ACTIVATION a)
@@ -34,21 +35,39 @@
    return RELU;
}

float activate(float x, ACTIVATION a){
float linear_activate(float x){return x;}
float sigmoid_activate(float x){return 1./(1. + exp(-x));}
float relu_activate(float x){return x*(x>0);}
float ramp_activate(float x){return x*(x>0)+.1*x;}
float tanh_activate(float x){return (exp(2*x)-1)/(exp(2*x)+1);}

float activate(float x, ACTIVATION a, float dropout)
{
    if((float)rand()/RAND_MAX < dropout) return 0;
    switch(a){
        case LINEAR:
            return x;
            return linear_activate(x)/(1-dropout);
        case SIGMOID:
            return 1./(1.+exp(-x));
            return sigmoid_activate(x)/(1-dropout);
        case RELU:
            return x*(x>0);
            return relu_activate(x)/(1-dropout);
        case RAMP:
            return x*(x>0) + .1*x;
            return ramp_activate(x)/(1-dropout);
        case TANH:
            return (exp(2*x)-1)/(exp(2*x)+1);
            return tanh_activate(x)/(1-dropout);
    }
    return 0;
}

void activate_array(float *x, const int n, const ACTIVATION a, float dropout)
{
    int i;
    for(i = 0; i < n; ++i){
        x[i] = activate(x[i], a, dropout);
    }
}


float gradient(float x, ACTIVATION a){
    switch(a){
        case LINEAR:
@@ -65,3 +84,48 @@
    return 0;
}

void gradient_array(const float *x, const int n, const ACTIVATION a, float *delta)
{
    int i;
    for(i = 0; i < n; ++i){
        delta[i] *= gradient(x[i], a);
    }
} 

#ifdef GPU

#include "opencl.h"
#include <math.h>

cl_kernel get_activation_kernel()
{
    static int init = 0;
    static cl_kernel kernel;
    if(!init){
        kernel = get_kernel("src/activations.cl", "activate_array", 0);
        init = 1;
    }
    return kernel;
}


void activate_array_ongpu(cl_mem x, int n, ACTIVATION a, float dropout) 
{
    cl_setup();
    cl_kernel kernel = get_activation_kernel();
    cl_command_queue queue = cl.queue;

    cl_uint i = 0;
    cl.error = clSetKernelArg(kernel, i++, sizeof(x), (void*) &x);
    cl.error = clSetKernelArg(kernel, i++, sizeof(n), (void*) &n);
    cl.error = clSetKernelArg(kernel, i++, sizeof(a), (void*) &a);
    cl.error = clSetKernelArg(kernel, i++, sizeof(dropout), 
        (void*) &dropout);
    check_error(cl);

    size_t gsize = n;

    clEnqueueNDRangeKernel(queue, kernel, 1, 0, &gsize, 0, 0, 0, 0);
    check_error(cl);
}
#endif