| | |
| | | #include "region_layer.h" |
| | | #include "activations.h" |
| | | #include "softmax_layer.h" |
| | | #include "blas.h" |
| | | #include "box.h" |
| | | #include "cuda.h" |
| | |
| | | #include <string.h> |
| | | #include <stdlib.h> |
| | | |
| | | region_layer make_region_layer(int batch, int inputs, int n, int side, int classes, int coords, int rescore) |
| | | #define DOABS 1 |
| | | |
| | | region_layer make_region_layer(int batch, int w, int h, int n, int classes, int coords, int max_boxes) |
| | | { |
| | | region_layer l = {0}; |
| | | l.type = REGION; |
| | | |
| | | l.n = n; |
| | | l.batch = batch; |
| | | l.inputs = inputs; |
| | | l.h = h; |
| | | l.w = w; |
| | | l.classes = classes; |
| | | l.coords = coords; |
| | | l.rescore = rescore; |
| | | l.side = side; |
| | | assert(side*side*((1 + l.coords)*l.n + l.classes) == inputs); |
| | | l.cost = calloc(1, sizeof(float)); |
| | | l.outputs = l.inputs; |
| | | l.truths = l.side*l.side*(1+l.coords+l.classes); |
| | | l.output = calloc(batch*l.outputs, sizeof(float)); |
| | | l.biases = calloc(n*2, sizeof(float)); |
| | | l.bias_updates = calloc(n*2, sizeof(float)); |
| | | l.outputs = h*w*n*(classes + coords + 1); |
| | | l.inputs = l.outputs; |
| | | l.max_boxes = max_boxes; |
| | | l.truths = max_boxes*(5); |
| | | l.delta = calloc(batch*l.outputs, sizeof(float)); |
| | | l.output = calloc(batch*l.outputs, sizeof(float)); |
| | | int i; |
| | | for(i = 0; i < n*2; ++i){ |
| | | l.biases[i] = .5; |
| | | } |
| | | |
| | | l.forward = forward_region_layer; |
| | | l.backward = backward_region_layer; |
| | | #ifdef GPU |
| | | l.forward_gpu = forward_region_layer_gpu; |
| | | l.backward_gpu = backward_region_layer_gpu; |
| | | l.output_gpu = cuda_make_array(l.output, batch*l.outputs); |
| | | l.delta_gpu = cuda_make_array(l.delta, batch*l.outputs); |
| | | #endif |
| | | |
| | | fprintf(stderr, "Region Layer\n"); |
| | | fprintf(stderr, "detection\n"); |
| | | srand(0); |
| | | |
| | | return l; |
| | | } |
| | | |
| | | void resize_region_layer(layer *l, int w, int h) |
| | | { |
| | | int old_w = l->w; |
| | | int old_h = l->h; |
| | | l->w = w; |
| | | l->h = h; |
| | | |
| | | l->outputs = h*w*l->n*(l->classes + l->coords + 1); |
| | | l->inputs = l->outputs; |
| | | |
| | | l->output = realloc(l->output, l->batch*l->outputs*sizeof(float)); |
| | | l->delta = realloc(l->delta, l->batch*l->outputs*sizeof(float)); |
| | | |
| | | #ifdef GPU |
| | | if (old_w < w || old_h < h) { |
| | | cuda_free(l->delta_gpu); |
| | | cuda_free(l->output_gpu); |
| | | |
| | | l->delta_gpu = cuda_make_array(l->delta, l->batch*l->outputs); |
| | | l->output_gpu = cuda_make_array(l->output, l->batch*l->outputs); |
| | | } |
| | | #endif |
| | | } |
| | | |
| | | box get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h) |
| | | { |
| | | box b; |
| | | b.x = (i + logistic_activate(x[index + 0])) / w; |
| | | b.y = (j + logistic_activate(x[index + 1])) / h; |
| | | b.w = exp(x[index + 2]) * biases[2*n]; |
| | | b.h = exp(x[index + 3]) * biases[2*n+1]; |
| | | if(DOABS){ |
| | | b.w = exp(x[index + 2]) * biases[2*n] / w; |
| | | b.h = exp(x[index + 3]) * biases[2*n+1] / h; |
| | | } |
| | | return b; |
| | | } |
| | | |
| | | float delta_region_box(box truth, float *x, float *biases, int n, int index, int i, int j, int w, int h, float *delta, float scale) |
| | | { |
| | | box pred = get_region_box(x, biases, n, index, i, j, w, h); |
| | | float iou = box_iou(pred, truth); |
| | | |
| | | float tx = (truth.x*w - i); |
| | | float ty = (truth.y*h - j); |
| | | float tw = log(truth.w / biases[2*n]); |
| | | float th = log(truth.h / biases[2*n + 1]); |
| | | if(DOABS){ |
| | | tw = log(truth.w*w / biases[2*n]); |
| | | th = log(truth.h*h / biases[2*n + 1]); |
| | | } |
| | | |
| | | delta[index + 0] = scale * (tx - logistic_activate(x[index + 0])) * logistic_gradient(logistic_activate(x[index + 0])); |
| | | delta[index + 1] = scale * (ty - logistic_activate(x[index + 1])) * logistic_gradient(logistic_activate(x[index + 1])); |
| | | delta[index + 2] = scale * (tw - x[index + 2]); |
| | | delta[index + 3] = scale * (th - x[index + 3]); |
| | | return iou; |
| | | } |
| | | |
| | | void delta_region_class(float *output, float *delta, int index, int class, int classes, tree *hier, float scale, float *avg_cat) |
| | | { |
| | | int i, n; |
| | | if(hier){ |
| | | float pred = 1; |
| | | while(class >= 0){ |
| | | pred *= output[index + class]; |
| | | int g = hier->group[class]; |
| | | int offset = hier->group_offset[g]; |
| | | for(i = 0; i < hier->group_size[g]; ++i){ |
| | | delta[index + offset + i] = scale * (0 - output[index + offset + i]); |
| | | } |
| | | delta[index + class] = scale * (1 - output[index + class]); |
| | | |
| | | class = hier->parent[class]; |
| | | } |
| | | *avg_cat += pred; |
| | | } else { |
| | | for(n = 0; n < classes; ++n){ |
| | | delta[index + n] = scale * (((n == class)?1 : 0) - output[index + n]); |
| | | if(n == class) *avg_cat += output[index + n]; |
| | | } |
| | | } |
| | | } |
| | | |
| | | float logit(float x) |
| | | { |
| | | return log(x/(1.-x)); |
| | | } |
| | | |
| | | float tisnan(float x) |
| | | { |
| | | return (x != x); |
| | | } |
| | | |
| | | void softmax_tree(float *input, int batch, int inputs, float temp, tree *hierarchy, float *output); |
| | | void forward_region_layer(const region_layer l, network_state state) |
| | | { |
| | | int locations = l.side*l.side; |
| | | int i,j; |
| | | int i,j,b,t,n; |
| | | int size = l.coords + l.classes + 1; |
| | | memcpy(l.output, state.input, l.outputs*l.batch*sizeof(float)); |
| | | int b; |
| | | if (l.softmax){ |
| | | for(b = 0; b < l.batch; ++b){ |
| | | int index = b*l.inputs; |
| | | for (i = 0; i < locations; ++i) { |
| | | int offset = i*l.classes; |
| | | softmax_array(l.output + index + offset, l.classes, |
| | | l.output + index + offset); |
| | | } |
| | | int offset = locations*l.classes; |
| | | activate_array(l.output + index + offset, locations*l.n*(1+l.coords), LOGISTIC); |
| | | } |
| | | } |
| | | if (l.object_logistic) { |
| | | for(b = 0; b < l.batch; ++b){ |
| | | int index = b*l.inputs; |
| | | int p_index = index + locations*l.classes; |
| | | activate_array(l.output + p_index, locations*l.n, LOGISTIC); |
| | | #ifndef GPU |
| | | flatten(l.output, l.w*l.h, size*l.n, l.batch, 1); |
| | | #endif |
| | | for (b = 0; b < l.batch; ++b){ |
| | | for(i = 0; i < l.h*l.w*l.n; ++i){ |
| | | int index = size*i + b*l.outputs; |
| | | l.output[index + 4] = logistic_activate(l.output[index + 4]); |
| | | } |
| | | } |
| | | |
| | | if (l.coord_logistic) { |
| | | for(b = 0; b < l.batch; ++b){ |
| | | int index = b*l.inputs; |
| | | int coord_index = index + locations*(l.classes + l.n); |
| | | activate_array(l.output + coord_index, locations*l.n*l.coords, LOGISTIC); |
| | | } |
| | | } |
| | | |
| | | if (l.class_logistic) { |
| | | for(b = 0; b < l.batch; ++b){ |
| | | int class_index = b*l.inputs; |
| | | activate_array(l.output + class_index, locations*l.classes, LOGISTIC); |
| | | } |
| | | } |
| | | |
| | | if(state.train){ |
| | | float avg_iou = 0; |
| | | float avg_cat = 0; |
| | | float avg_allcat = 0; |
| | | float avg_obj = 0; |
| | | float avg_anyobj = 0; |
| | | int count = 0; |
| | | *(l.cost) = 0; |
| | | int size = l.inputs * l.batch; |
| | | memset(l.delta, 0, size * sizeof(float)); |
| | | #ifndef GPU |
| | | if (l.softmax_tree){ |
| | | for (b = 0; b < l.batch; ++b){ |
| | | int index = b*l.inputs; |
| | | for (i = 0; i < locations; ++i) { |
| | | int truth_index = (b*locations + i)*(1+l.coords+l.classes); |
| | | int is_obj = state.truth[truth_index]; |
| | | for (j = 0; j < l.n; ++j) { |
| | | int p_index = index + locations*l.classes + i*l.n + j; |
| | | l.delta[p_index] = l.noobject_scale*(0 - l.output[p_index]); |
| | | *(l.cost) += l.noobject_scale*pow(l.output[p_index], 2); |
| | | avg_anyobj += l.output[p_index]; |
| | | } |
| | | |
| | | int best_index = -1; |
| | | float best_iou = 0; |
| | | float best_rmse = 20; |
| | | |
| | | if (!is_obj){ |
| | | continue; |
| | | } |
| | | |
| | | int class_index = index + i*l.classes; |
| | | for(j = 0; j < l.classes; ++j) { |
| | | l.delta[class_index+j] = l.class_scale * (state.truth[truth_index+1+j] - l.output[class_index+j]); |
| | | *(l.cost) += l.class_scale * pow(state.truth[truth_index+1+j] - l.output[class_index+j], 2); |
| | | if(state.truth[truth_index + 1 + j]) avg_cat += l.output[class_index+j]; |
| | | avg_allcat += l.output[class_index+j]; |
| | | } |
| | | |
| | | box truth = float_to_box(state.truth + truth_index + 1 + l.classes); |
| | | truth.x /= l.side; |
| | | truth.y /= l.side; |
| | | |
| | | for(j = 0; j < l.n; ++j){ |
| | | int box_index = index + locations*(l.classes + l.n) + (i*l.n + j) * l.coords; |
| | | box out = float_to_box(l.output + box_index); |
| | | out.x /= l.side; |
| | | out.y /= l.side; |
| | | |
| | | if (l.sqrt){ |
| | | out.w = out.w*out.w; |
| | | out.h = out.h*out.h; |
| | | } |
| | | |
| | | float iou = box_iou(out, truth); |
| | | //iou = 0; |
| | | float rmse = box_rmse(out, truth); |
| | | if(best_iou > 0 || iou > 0){ |
| | | if(iou > best_iou){ |
| | | best_iou = iou; |
| | | best_index = j; |
| | | } |
| | | }else{ |
| | | if(rmse < best_rmse){ |
| | | best_rmse = rmse; |
| | | best_index = j; |
| | | } |
| | | } |
| | | } |
| | | |
| | | if(l.forced){ |
| | | if(truth.w*truth.h < .1){ |
| | | best_index = 1; |
| | | }else{ |
| | | best_index = 0; |
| | | } |
| | | } |
| | | |
| | | int box_index = index + locations*(l.classes + l.n) + (i*l.n + best_index) * l.coords; |
| | | int tbox_index = truth_index + 1 + l.classes; |
| | | |
| | | box out = float_to_box(l.output + box_index); |
| | | out.x /= l.side; |
| | | out.y /= l.side; |
| | | if (l.sqrt) { |
| | | out.w = out.w*out.w; |
| | | out.h = out.h*out.h; |
| | | } |
| | | float iou = box_iou(out, truth); |
| | | |
| | | //printf("%d", best_index); |
| | | int p_index = index + locations*l.classes + i*l.n + best_index; |
| | | *(l.cost) -= l.noobject_scale * pow(l.output[p_index], 2); |
| | | *(l.cost) += l.object_scale * pow(1-l.output[p_index], 2); |
| | | avg_obj += l.output[p_index]; |
| | | l.delta[p_index] = l.object_scale * (1.-l.output[p_index]); |
| | | |
| | | if(l.rescore){ |
| | | l.delta[p_index] = l.object_scale * (iou - l.output[p_index]); |
| | | } |
| | | |
| | | l.delta[box_index+0] = l.coord_scale*(state.truth[tbox_index + 0] - l.output[box_index + 0]); |
| | | l.delta[box_index+1] = l.coord_scale*(state.truth[tbox_index + 1] - l.output[box_index + 1]); |
| | | l.delta[box_index+2] = l.coord_scale*(state.truth[tbox_index + 2] - l.output[box_index + 2]); |
| | | l.delta[box_index+3] = l.coord_scale*(state.truth[tbox_index + 3] - l.output[box_index + 3]); |
| | | if(l.sqrt){ |
| | | l.delta[box_index+2] = l.coord_scale*(sqrt(state.truth[tbox_index + 2]) - l.output[box_index + 2]); |
| | | l.delta[box_index+3] = l.coord_scale*(sqrt(state.truth[tbox_index + 3]) - l.output[box_index + 3]); |
| | | } |
| | | |
| | | *(l.cost) += pow(1-iou, 2); |
| | | avg_iou += iou; |
| | | ++count; |
| | | for(i = 0; i < l.h*l.w*l.n; ++i){ |
| | | int index = size*i + b*l.outputs; |
| | | softmax_tree(l.output + index + 5, 1, 0, 1, l.softmax_tree, l.output + index + 5); |
| | | } |
| | | if(l.softmax){ |
| | | gradient_array(l.output + index + locations*l.classes, locations*l.n*(1+l.coords), |
| | | LOGISTIC, l.delta + index + locations*l.classes); |
| | | } |
| | | if (l.object_logistic) { |
| | | int p_index = index + locations*l.classes; |
| | | gradient_array(l.output + p_index, locations*l.n, LOGISTIC, l.delta + p_index); |
| | | } |
| | | |
| | | if (l.class_logistic) { |
| | | int class_index = index; |
| | | gradient_array(l.output + class_index, locations*l.classes, LOGISTIC, l.delta + class_index); |
| | | } |
| | | |
| | | if (l.coord_logistic) { |
| | | int coord_index = index + locations*(l.classes + l.n); |
| | | gradient_array(l.output + coord_index, locations*l.n*l.coords, LOGISTIC, l.delta + coord_index); |
| | | } |
| | | //printf("\n"); |
| | | } |
| | | printf("Region Avg IOU: %f, Pos Cat: %f, All Cat: %f, Pos Obj: %f, Any Obj: %f, count: %d\n", avg_iou/count, avg_cat/count, avg_allcat/(count*l.classes), avg_obj/count, avg_anyobj/(l.batch*locations*l.n), count); |
| | | } else if (l.softmax){ |
| | | for (b = 0; b < l.batch; ++b){ |
| | | for(i = 0; i < l.h*l.w*l.n; ++i){ |
| | | int index = size*i + b*l.outputs; |
| | | softmax(l.output + index + 5, l.classes, 1, l.output + index + 5, 1); |
| | | } |
| | | } |
| | | } |
| | | #endif |
| | | if(!state.train) return; |
| | | memset(l.delta, 0, l.outputs * l.batch * sizeof(float)); |
| | | float avg_iou = 0; |
| | | float recall = 0; |
| | | float avg_cat = 0; |
| | | float avg_obj = 0; |
| | | float avg_anyobj = 0; |
| | | int count = 0; |
| | | int class_count = 0; |
| | | *(l.cost) = 0; |
| | | for (b = 0; b < l.batch; ++b) { |
| | | if(l.softmax_tree){ |
| | | int onlyclass = 0; |
| | | for(t = 0; t < l.max_boxes; ++t){ |
| | | box truth = float_to_box(state.truth + t*5 + b*l.truths); |
| | | if(!truth.x) break; |
| | | int class = state.truth[t*5 + b*l.truths + 4]; |
| | | float maxp = 0; |
| | | int maxi = 0; |
| | | if(truth.x > 100000 && truth.y > 100000){ |
| | | for(n = 0; n < l.n*l.w*l.h; ++n){ |
| | | int index = size*n + b*l.outputs + 5; |
| | | float scale = l.output[index-1]; |
| | | float p = scale*get_hierarchy_probability(l.output + index, l.softmax_tree, class); |
| | | if(p > maxp){ |
| | | maxp = p; |
| | | maxi = n; |
| | | } |
| | | } |
| | | int index = size*maxi + b*l.outputs + 5; |
| | | delta_region_class(l.output, l.delta, index, class, l.classes, l.softmax_tree, l.class_scale, &avg_cat); |
| | | ++class_count; |
| | | onlyclass = 1; |
| | | break; |
| | | } |
| | | } |
| | | if(onlyclass) continue; |
| | | } |
| | | for (j = 0; j < l.h; ++j) { |
| | | for (i = 0; i < l.w; ++i) { |
| | | for (n = 0; n < l.n; ++n) { |
| | | int index = size*(j*l.w*l.n + i*l.n + n) + b*l.outputs; |
| | | box pred = get_region_box(l.output, l.biases, n, index, i, j, l.w, l.h); |
| | | float best_iou = 0; |
| | | int best_class = -1; |
| | | for(t = 0; t < l.max_boxes; ++t){ |
| | | box truth = float_to_box(state.truth + t*5 + b*l.truths); |
| | | if(!truth.x) break; |
| | | float iou = box_iou(pred, truth); |
| | | if (iou > best_iou) { |
| | | best_class = state.truth[t*5 + b*l.truths + 4]; |
| | | best_iou = iou; |
| | | } |
| | | } |
| | | avg_anyobj += l.output[index + 4]; |
| | | l.delta[index + 4] = l.noobject_scale * ((0 - l.output[index + 4]) * logistic_gradient(l.output[index + 4])); |
| | | if(l.classfix == -1) l.delta[index + 4] = l.noobject_scale * ((best_iou - l.output[index + 4]) * logistic_gradient(l.output[index + 4])); |
| | | else{ |
| | | if (best_iou > l.thresh) { |
| | | l.delta[index + 4] = 0; |
| | | if(l.classfix > 0){ |
| | | delta_region_class(l.output, l.delta, index + 5, best_class, l.classes, l.softmax_tree, l.class_scale*(l.classfix == 2 ? l.output[index + 4] : 1), &avg_cat); |
| | | ++class_count; |
| | | } |
| | | } |
| | | } |
| | | |
| | | if(*(state.net.seen) < 12800){ |
| | | box truth = {0}; |
| | | truth.x = (i + .5)/l.w; |
| | | truth.y = (j + .5)/l.h; |
| | | truth.w = l.biases[2*n]; |
| | | truth.h = l.biases[2*n+1]; |
| | | if(DOABS){ |
| | | truth.w = l.biases[2*n]/l.w; |
| | | truth.h = l.biases[2*n+1]/l.h; |
| | | } |
| | | delta_region_box(truth, l.output, l.biases, n, index, i, j, l.w, l.h, l.delta, .01); |
| | | } |
| | | } |
| | | } |
| | | } |
| | | for(t = 0; t < l.max_boxes; ++t){ |
| | | box truth = float_to_box(state.truth + t*5 + b*l.truths); |
| | | |
| | | if(!truth.x) break; |
| | | float best_iou = 0; |
| | | int best_index = 0; |
| | | int best_n = 0; |
| | | i = (truth.x * l.w); |
| | | j = (truth.y * l.h); |
| | | //printf("%d %f %d %f\n", i, truth.x*l.w, j, truth.y*l.h); |
| | | box truth_shift = truth; |
| | | truth_shift.x = 0; |
| | | truth_shift.y = 0; |
| | | //printf("index %d %d\n",i, j); |
| | | for(n = 0; n < l.n; ++n){ |
| | | int index = size*(j*l.w*l.n + i*l.n + n) + b*l.outputs; |
| | | box pred = get_region_box(l.output, l.biases, n, index, i, j, l.w, l.h); |
| | | if(l.bias_match){ |
| | | pred.w = l.biases[2*n]; |
| | | pred.h = l.biases[2*n+1]; |
| | | if(DOABS){ |
| | | pred.w = l.biases[2*n]/l.w; |
| | | pred.h = l.biases[2*n+1]/l.h; |
| | | } |
| | | } |
| | | //printf("pred: (%f, %f) %f x %f\n", pred.x, pred.y, pred.w, pred.h); |
| | | pred.x = 0; |
| | | pred.y = 0; |
| | | float iou = box_iou(pred, truth_shift); |
| | | if (iou > best_iou){ |
| | | best_index = index; |
| | | best_iou = iou; |
| | | best_n = n; |
| | | } |
| | | } |
| | | //printf("%d %f (%f, %f) %f x %f\n", best_n, best_iou, truth.x, truth.y, truth.w, truth.h); |
| | | |
| | | float iou = delta_region_box(truth, l.output, l.biases, best_n, best_index, i, j, l.w, l.h, l.delta, l.coord_scale); |
| | | if(iou > .5) recall += 1; |
| | | avg_iou += iou; |
| | | |
| | | //l.delta[best_index + 4] = iou - l.output[best_index + 4]; |
| | | avg_obj += l.output[best_index + 4]; |
| | | l.delta[best_index + 4] = l.object_scale * (1 - l.output[best_index + 4]) * logistic_gradient(l.output[best_index + 4]); |
| | | if (l.rescore) { |
| | | l.delta[best_index + 4] = l.object_scale * (iou - l.output[best_index + 4]) * logistic_gradient(l.output[best_index + 4]); |
| | | } |
| | | |
| | | |
| | | int class = state.truth[t*5 + b*l.truths + 4]; |
| | | if (l.map) class = l.map[class]; |
| | | delta_region_class(l.output, l.delta, best_index + 5, class, l.classes, l.softmax_tree, l.class_scale, &avg_cat); |
| | | ++count; |
| | | ++class_count; |
| | | } |
| | | } |
| | | //printf("\n"); |
| | | #ifndef GPU |
| | | flatten(l.delta, l.w*l.h, size*l.n, l.batch, 0); |
| | | #endif |
| | | *(l.cost) = pow(mag_array(l.delta, l.outputs * l.batch), 2); |
| | | printf("Region Avg IOU: %f, Class: %f, Obj: %f, No Obj: %f, Avg Recall: %f, count: %d\n", avg_iou/count, avg_cat/class_count, avg_obj/count, avg_anyobj/(l.w*l.h*l.n*l.batch), recall/count, count); |
| | | } |
| | | |
| | | void backward_region_layer(const region_layer l, network_state state) |
| | |
| | | axpy_cpu(l.batch*l.inputs, 1, l.delta, 1, state.delta, 1); |
| | | } |
| | | |
| | | void get_region_boxes(layer l, int w, int h, float thresh, float **probs, box *boxes, int only_objectness, int *map) |
| | | { |
| | | int i,j,n; |
| | | float *predictions = l.output; |
| | | for (i = 0; i < l.w*l.h; ++i){ |
| | | int row = i / l.w; |
| | | int col = i % l.w; |
| | | for(n = 0; n < l.n; ++n){ |
| | | int index = i*l.n + n; |
| | | int p_index = index * (l.classes + 5) + 4; |
| | | float scale = predictions[p_index]; |
| | | if(l.classfix == -1 && scale < .5) scale = 0; |
| | | int box_index = index * (l.classes + 5); |
| | | boxes[index] = get_region_box(predictions, l.biases, n, box_index, col, row, l.w, l.h); |
| | | boxes[index].x *= w; |
| | | boxes[index].y *= h; |
| | | boxes[index].w *= w; |
| | | boxes[index].h *= h; |
| | | |
| | | int class_index = index * (l.classes + 5) + 5; |
| | | if(l.softmax_tree){ |
| | | |
| | | hierarchy_predictions(predictions + class_index, l.classes, l.softmax_tree, 0); |
| | | int found = 0; |
| | | if(map){ |
| | | for(j = 0; j < 200; ++j){ |
| | | float prob = scale*predictions[class_index+map[j]]; |
| | | probs[index][j] = (prob > thresh) ? prob : 0; |
| | | } |
| | | } else { |
| | | for(j = l.classes - 1; j >= 0; --j){ |
| | | if(!found && predictions[class_index + j] > .5){ |
| | | found = 1; |
| | | } else { |
| | | predictions[class_index + j] = 0; |
| | | } |
| | | float prob = predictions[class_index+j]; |
| | | probs[index][j] = (scale > thresh) ? prob : 0; |
| | | } |
| | | } |
| | | } else { |
| | | for(j = 0; j < l.classes; ++j){ |
| | | float prob = scale*predictions[class_index+j]; |
| | | probs[index][j] = (prob > thresh) ? prob : 0; |
| | | } |
| | | } |
| | | if(only_objectness){ |
| | | probs[index][0] = scale; |
| | | } |
| | | } |
| | | } |
| | | } |
| | | |
| | | #ifdef GPU |
| | | |
| | | void forward_region_layer_gpu(const region_layer l, network_state state) |
| | | { |
| | | /* |
| | | if(!state.train){ |
| | | copy_ongpu(l.batch*l.inputs, state.input, 1, l.output_gpu, 1); |
| | | return; |
| | | } |
| | | */ |
| | | flatten_ongpu(state.input, l.h*l.w, l.n*(l.coords + l.classes + 1), l.batch, 1, l.output_gpu); |
| | | if(l.softmax_tree){ |
| | | int i; |
| | | int count = 5; |
| | | for (i = 0; i < l.softmax_tree->groups; ++i) { |
| | | int group_size = l.softmax_tree->group_size[i]; |
| | | softmax_gpu(l.output_gpu+count, group_size, l.classes + 5, l.w*l.h*l.n*l.batch, 1, l.output_gpu + count); |
| | | count += group_size; |
| | | } |
| | | }else if (l.softmax){ |
| | | softmax_gpu(l.output_gpu+5, l.classes, l.classes + 5, l.w*l.h*l.n*l.batch, 1, l.output_gpu + 5); |
| | | } |
| | | |
| | | float *in_cpu = calloc(l.batch*l.inputs, sizeof(float)); |
| | | float *truth_cpu = 0; |
| | | if(state.truth){ |
| | | int num_truth = l.batch*l.side*l.side*(1+l.coords+l.classes); |
| | | int num_truth = l.batch*l.truths; |
| | | truth_cpu = calloc(num_truth, sizeof(float)); |
| | | cuda_pull_array(state.truth, truth_cpu, num_truth); |
| | | } |
| | | cuda_pull_array(state.input, in_cpu, l.batch*l.inputs); |
| | | network_state cpu_state; |
| | | cuda_pull_array(l.output_gpu, in_cpu, l.batch*l.inputs); |
| | | cudaStreamSynchronize(get_cuda_stream()); |
| | | network_state cpu_state = state; |
| | | cpu_state.train = state.train; |
| | | cpu_state.truth = truth_cpu; |
| | | cpu_state.input = in_cpu; |
| | | forward_region_layer(l, cpu_state); |
| | | cuda_push_array(l.output_gpu, l.output, l.batch*l.outputs); |
| | | cuda_push_array(l.delta_gpu, l.delta, l.batch*l.inputs); |
| | | //cuda_push_array(l.output_gpu, l.output, l.batch*l.outputs); |
| | | free(cpu_state.input); |
| | | if(!state.train) return; |
| | | cuda_push_array(l.delta_gpu, l.delta, l.batch*l.outputs); |
| | | cudaStreamSynchronize(get_cuda_stream()); |
| | | if(cpu_state.truth) free(cpu_state.truth); |
| | | } |
| | | |
| | | void backward_region_layer_gpu(region_layer l, network_state state) |
| | | { |
| | | axpy_ongpu(l.batch*l.inputs, 1, l.delta_gpu, 1, state.delta, 1); |
| | | //copy_ongpu(l.batch*l.inputs, l.delta_gpu, 1, state.delta, 1); |
| | | flatten_ongpu(l.delta_gpu, l.h*l.w, l.n*(l.coords + l.classes + 1), l.batch, 0, state.delta); |
| | | } |
| | | #endif |
| | | |