Detection is back, baby\!

Joseph Redmon

2015-02-11 0f645836f193e75c4c3b718369e6fab15b5d19c5

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
extern "C" {
#include "softmax_layer.h"
#include "cuda.h"
#include "blas.h"
}
 
#define BLOCK 256
 
__global__ void forward_softmax_layer_kernel(int n, int batch, float *input, float *output)
{
    int b = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
    if(b >= batch) return;
 
    int i;
    float sum = 0;
    float largest = -INFINITY;
    for(i = 0; i < n; ++i){
        int val = input[i+b*n];
        largest = (val>largest) ? val : largest;
    }
    for(i = 0; i < n; ++i){
        sum += exp(input[i+b*n]-largest);
    }
    sum = (sum != 0) ? largest+log(sum) : largest-100;
    for(i = 0; i < n; ++i){
        output[i+b*n] = exp(input[i+b*n]-sum);
    }
}
 
extern "C" void pull_softmax_layer_output(const softmax_layer layer)
{
    cuda_pull_array(layer.output_gpu, layer.output, layer.inputs*layer.batch);
}
 
extern "C" void forward_softmax_layer_gpu(const softmax_layer layer, float *input)
{
    forward_softmax_layer_kernel<<<cuda_gridsize(layer.batch), BLOCK>>>(layer.inputs, layer.batch, input, layer.output_gpu);
    check_error(cudaPeekAtLastError());
 
    /*
    cl_read_array(layer.output_cl, layer.output, layer.inputs*layer.batch);
    int z;
    for(z = 0; z < layer.inputs*layer.batch; ++z) printf("%f,",layer.output[z]);
    */
}
 
extern "C" void backward_softmax_layer_gpu(const softmax_layer layer, float *delta)
{
    copy_ongpu(layer.batch*layer.inputs, layer.delta_gpu, 1, delta, 1);
}
 
/* This is if you want softmax w/o log-loss classification. You probably don't.
   int i,j,b;
   for(b = 0; b < layer.batch; ++b){
   for(i = 0; i < layer.inputs; ++i){
   for(j = 0; j < layer.inputs; ++j){
   int d = (i==j);
   layer.jacobian[b*layer.inputs*layer.inputs + i*layer.inputs + j] = 
   layer.output[b*layer.inputs + i] * (d - layer.output[b*layer.inputs + j]);
   }
   }
   }
   for(b = 0; b < layer.batch; ++b){
   int M = layer.inputs;
   int N = 1;
   int K = layer.inputs;
   float *A = layer.jacobian + b*layer.inputs*layer.inputs;
   float *B = layer.delta + b*layer.inputs;
   float *C = delta + b*layer.inputs;
   gemm(0,0,M,N,K,1,A,K,B,N,0,C,N);
   }
 */