From d8adaf8ea6a31a380f6bf1fe65e88b661d3bb51e Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Fri, 21 Oct 2016 20:16:43 +0000
Subject: [PATCH] tree stuff
---
src/softmax_layer.c | 117 +++++++++++++++++++++++++++++++++++++++++++++-------------
1 files changed, 90 insertions(+), 27 deletions(-)
diff --git a/src/softmax_layer.c b/src/softmax_layer.c
index aa5ab06..2a34cae 100644
--- a/src/softmax_layer.c
+++ b/src/softmax_layer.c
@@ -5,48 +5,111 @@
#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)
{
+ assert(inputs%groups == 0);
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));
+ 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_gpu = cuda_make_array(layer->output, inputs*batch);
- layer->delta_gpu = cuda_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 forward_softmax_layer(const softmax_layer l, network_state state)
{
- 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];
+ int b;
+ int inputs = l.inputs / l.groups;
+ int batch = l.batch * l.groups;
+ if(l.softmax_tree){
+ for(b = 0; b < batch; ++b){
+ int i;
+ int count = 0;
+ for(i = 0; i < l.softmax_tree->groups; ++i){
+ int group_size = l.softmax_tree->group_size[i];
+ softmax(state.input+b*inputs + count, group_size, l.temperature, l.output+b*inputs + count);
+ count += group_size;
+ }
}
- 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);
+ } else {
+ for(b = 0; b < batch; ++b){
+ softmax(state.input+b*inputs, inputs, l.temperature, l.output+b*inputs);
}
}
}
-void backward_softmax_layer(const softmax_layer layer, float *delta)
+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];
}
}
+#ifdef GPU
+
+void pull_softmax_layer_output(const softmax_layer layer)
+{
+ cuda_pull_array(layer.output_gpu, layer.output, layer.inputs*layer.batch);
+}
+
+void forward_softmax_layer_gpu(const softmax_layer l, network_state state)
+{
+ int inputs = l.inputs / l.groups;
+ int batch = l.batch * l.groups;
+ int b;
+ if(l.softmax_tree){
+ if(0){
+ float *buff = calloc(inputs * batch, sizeof(float));
+ cuda_pull_array(state.input, buff, batch * inputs);
+ state.input = buff;
+ forward_softmax_layer(l, state);
+ cuda_push_array(l.output_gpu, l.output, batch*inputs);
+ free(buff);
+ } else {
+ int i;
+ const int nstreams = 32;
+ cudaStream_t streams[nstreams];
+ for (i = 0; i < nstreams; ++i) {
+ cudaStreamCreate(&streams[i]);
+ }
+ for (b = 0; b < batch; ++b) {
+ 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+b*inputs + count, group_size, 1, l.temperature, l.output_gpu+b*inputs + count, streams[(b*l.softmax_tree->groups + i) % nstreams]);
+ count += group_size;
+ }
+ }
+ for(i = 0; i < nstreams; ++i){
+ cudaStreamDestroy(streams[i]);
+ }
+ }
+ } else {
+ softmax_gpu(state.input, inputs, batch, l.temperature, l.output_gpu, 0);
+ }
+}
+
+void backward_softmax_layer_gpu(const softmax_layer layer, network_state state)
+{
+ axpy_ongpu(layer.batch*layer.inputs, 1, layer.delta_gpu, 1, state.delta, 1);
+}
+
+#endif
--
Gitblit v1.10.0