From 91f95c715bff84094fc18bad6a8f938291b9b0f5 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Mon, 24 Oct 2016 20:32:49 +0000
Subject: [PATCH] tree things, tree stuff
---
src/softmax_layer.c | 35 +++++++----------------------------
1 files changed, 7 insertions(+), 28 deletions(-)
diff --git a/src/softmax_layer.c b/src/softmax_layer.c
index 2a34cae..31f3e03 100644
--- a/src/softmax_layer.c
+++ b/src/softmax_layer.c
@@ -73,37 +73,16 @@
{
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]);
- }
+ 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, batch, l.temperature, l.output_gpu, 0);
+ softmax_gpu(state.input, inputs, inputs, batch, l.temperature, l.output_gpu);
}
}
--
Gitblit v1.10.0