From 028696bf15efeca3acb3db8c42a96f7b9e0f55ff Mon Sep 17 00:00:00 2001
From: iovodov <b@ovdv.ru>
Date: Thu, 03 May 2018 13:33:46 +0000
Subject: [PATCH] Output improvements for detector results: When printing detector results, output was done in random order, obfuscating results for interpreting. Now: 1. Text output includes coordinates of rects in (left,right,top,bottom in pixels) along with label and score 2. Text output is sorted by rect lefts to simplify finding appropriate rects on image 3. If several class probs are > thresh for some detection, the most probable is written first and coordinates for others are not repeated 4. Rects are imprinted in image in order by their best class prob, so most probable rects are always on top and not overlayed by less probable ones 5. Most probable label for rect is always written first Also: 6. Message about low GPU memory include required amount
---
src/softmax_layer.c | 149 +++++++++++++++++++++----------------------------
1 files changed, 65 insertions(+), 84 deletions(-)
diff --git a/src/softmax_layer.c b/src/softmax_layer.c
index ffc028f..27f73fd 100644
--- a/src/softmax_layer.c
+++ b/src/softmax_layer.c
@@ -1,51 +1,70 @@
#include "softmax_layer.h"
-#include "mini_blas.h"
+#include "blas.h"
+#include "cuda.h"
#include <float.h>
#include <math.h>
#include <stdlib.h>
#include <stdio.h>
+#include <assert.h>
-softmax_layer *make_softmax_layer(int batch, int inputs)
+softmax_layer make_softmax_layer(int batch, int inputs, int groups)
{
- fprintf(stderr, "Softmax Layer: %d inputs\n", inputs);
- softmax_layer *layer = calloc(1, sizeof(softmax_layer));
- layer->batch = batch;
- layer->inputs = inputs;
- layer->output = calloc(inputs*batch, sizeof(float));
- layer->delta = calloc(inputs*batch, sizeof(float));
- layer->jacobian = calloc(inputs*inputs*batch, sizeof(float));
+ assert(inputs%groups == 0);
+ fprintf(stderr, "softmax %4d\n", inputs);
+ softmax_layer l = {0};
+ l.type = SOFTMAX;
+ l.batch = batch;
+ l.groups = groups;
+ l.inputs = inputs;
+ l.outputs = inputs;
+ l.output = calloc(inputs*batch, sizeof(float));
+ l.delta = calloc(inputs*batch, sizeof(float));
+
+ l.forward = forward_softmax_layer;
+ l.backward = backward_softmax_layer;
#ifdef GPU
- layer->output_cl = cl_make_array(layer->output, inputs*batch);
- layer->delta_cl = cl_make_array(layer->delta, inputs*batch);
+ l.forward_gpu = forward_softmax_layer_gpu;
+ l.backward_gpu = backward_softmax_layer_gpu;
+
+ l.output_gpu = cuda_make_array(l.output, inputs*batch);
+ l.delta_gpu = cuda_make_array(l.delta, inputs*batch);
#endif
- return layer;
+ return l;
}
-void forward_softmax_layer(const softmax_layer layer, float *input)
+void softmax_tree(float *input, int batch, int inputs, float temp, tree *hierarchy, float *output)
{
- int i,b;
- for(b = 0; b < layer.batch; ++b){
- float sum = 0;
- float largest = -FLT_MAX;
- for(i = 0; i < layer.inputs; ++i){
- if(input[i+b*layer.inputs] > largest) largest = input[i+b*layer.inputs];
- }
- for(i = 0; i < layer.inputs; ++i){
- sum += exp(input[i+b*layer.inputs]-largest);
- }
- if(sum) sum = largest+log(sum);
- else sum = largest-100;
- for(i = 0; i < layer.inputs; ++i){
- layer.output[i+b*layer.inputs] = exp(input[i+b*layer.inputs]-sum);
+ int b;
+ for(b = 0; b < batch; ++b){
+ int i;
+ int count = 0;
+ for(i = 0; i < hierarchy->groups; ++i){
+ int group_size = hierarchy->group_size[i];
+ softmax(input+b*inputs + count, group_size, temp, output+b*inputs + count, 1);
+ count += group_size;
}
}
}
-void backward_softmax_layer(const softmax_layer layer, float *delta)
+void forward_softmax_layer(const softmax_layer l, network_state state)
+{
+ int b;
+ int inputs = l.inputs / l.groups;
+ int batch = l.batch * l.groups;
+ if(l.softmax_tree){
+ softmax_tree(state.input, batch, inputs, l.temperature, l.softmax_tree, l.output);
+ } else {
+ for(b = 0; b < batch; ++b){
+ softmax(state.input+b*inputs, inputs, l.temperature, l.output+b*inputs, 1);
+ }
+ }
+}
+
+void backward_softmax_layer(const softmax_layer l, network_state state)
{
int i;
- for(i = 0; i < layer.inputs*layer.batch; ++i){
- delta[i] = layer.delta[i];
+ for(i = 0; i < l.inputs*l.batch; ++i){
+ state.delta[i] += l.delta[i];
}
}
@@ -53,67 +72,29 @@
void pull_softmax_layer_output(const softmax_layer layer)
{
- cl_read_array(layer.output_cl, layer.output, layer.inputs*layer.batch);
+ cuda_pull_array(layer.output_gpu, layer.output, layer.inputs*layer.batch);
}
-cl_kernel get_softmax_forward_kernel()
+void forward_softmax_layer_gpu(const softmax_layer l, network_state state)
{
- static int init = 0;
- static cl_kernel kernel;
- if(!init){
- kernel = get_kernel("src/softmax_layer.cl", "forward", 0);
- init = 1;
+ int inputs = l.inputs / l.groups;
+ int batch = l.batch * l.groups;
+ if(l.softmax_tree){
+ int i;
+ int count = 0;
+ for (i = 0; i < l.softmax_tree->groups; ++i) {
+ int group_size = l.softmax_tree->group_size[i];
+ softmax_gpu(state.input+count, group_size, inputs, batch, l.temperature, l.output_gpu + count);
+ count += group_size;
+ }
+ } else {
+ softmax_gpu(state.input, inputs, inputs, batch, l.temperature, l.output_gpu);
}
- return kernel;
}
-void forward_softmax_layer_gpu(const softmax_layer layer, cl_mem input)
+void backward_softmax_layer_gpu(const softmax_layer layer, network_state state)
{
- cl_kernel kernel = get_softmax_forward_kernel();
- cl_command_queue queue = cl.queue;
-
- cl_uint i = 0;
- cl.error = clSetKernelArg(kernel, i++, sizeof(layer.inputs), (void*) &layer.inputs);
- cl.error = clSetKernelArg(kernel, i++, sizeof(input), (void*) &input);
- cl.error = clSetKernelArg(kernel, i++, sizeof(layer.output_cl), (void*) &layer.output_cl);
- check_error(cl);
-
- const size_t global_size[] = {layer.batch};
-
- cl.error = clEnqueueNDRangeKernel(queue, kernel, 1, 0, global_size, 0, 0, 0, 0);
- check_error(cl);
-
- /*
- 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]);
- */
+ axpy_ongpu(layer.batch*layer.inputs, 1, layer.delta_gpu, 1, state.delta, 1);
}
-void backward_softmax_layer_gpu(const softmax_layer layer, cl_mem delta)
-{
- copy_ongpu(layer.batch*layer.inputs, layer.delta_cl, 1, delta, 1);
-}
#endif
-
-/* 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);
- }
- */
--
Gitblit v1.10.0