From 989ab8c38a02fa7ea9c25108151736c62e81c972 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Fri, 24 Apr 2015 17:27:50 +0000
Subject: [PATCH] IOU loss function
---
src/crop_layer.c | 68 +++++++++------------------------
1 files changed, 19 insertions(+), 49 deletions(-)
diff --git a/src/crop_layer.c b/src/crop_layer.c
index 2a5007a..7ae4aa5 100644
--- a/src/crop_layer.c
+++ b/src/crop_layer.c
@@ -1,4 +1,5 @@
#include "crop_layer.h"
+#include "cuda.h"
#include <stdio.h>
image get_crop_image(crop_layer layer)
@@ -6,10 +7,10 @@
int h = layer.crop_height;
int w = layer.crop_width;
int c = layer.c;
- return float_to_image(h,w,c,layer.output);
+ return float_to_image(w,h,c,layer.output);
}
-crop_layer *make_crop_layer(int batch, int h, int w, int c, int crop_height, int crop_width, int flip)
+crop_layer *make_crop_layer(int batch, int h, int w, int c, int crop_height, int crop_width, int flip, float angle, float saturation, float exposure)
{
fprintf(stderr, "Crop Layer: %d x %d -> %d x %d x %d image\n", h,w,crop_height,crop_width,c);
crop_layer *layer = calloc(1, sizeof(crop_layer));
@@ -18,23 +19,34 @@
layer->w = w;
layer->c = c;
layer->flip = flip;
+ layer->angle = angle;
+ layer->saturation = saturation;
+ layer->exposure = exposure;
layer->crop_width = crop_width;
layer->crop_height = crop_height;
layer->output = calloc(crop_width*crop_height * c*batch, sizeof(float));
#ifdef GPU
- layer->output_cl = cl_make_array(layer->output, crop_width*crop_height*c*batch);
+ layer->output_gpu = cuda_make_array(layer->output, crop_width*crop_height*c*batch);
+ layer->rand_gpu = cuda_make_array(0, layer->batch*8);
#endif
return layer;
}
-void forward_crop_layer(const crop_layer layer, float *input)
+void forward_crop_layer(const crop_layer layer, network_state state)
{
int i,j,c,b,row,col;
int index;
int count = 0;
int flip = (layer.flip && rand()%2);
- int dh = rand()%(layer.h - layer.crop_height);
- int dw = rand()%(layer.w - layer.crop_width);
+ int dh = rand()%(layer.h - layer.crop_height + 1);
+ int dw = rand()%(layer.w - layer.crop_width + 1);
+ float scale = 2;
+ float trans = -1;
+ if(!state.train){
+ flip = 0;
+ dh = (layer.h - layer.crop_height)/2;
+ dw = (layer.w - layer.crop_width)/2;
+ }
for(b = 0; b < layer.batch; ++b){
for(c = 0; c < layer.c; ++c){
for(i = 0; i < layer.crop_height; ++i){
@@ -46,52 +58,10 @@
}
row = i + dh;
index = col+layer.w*(row+layer.h*(c + layer.c*b));
- layer.output[count++] = input[index];
+ layer.output[count++] = state.input[index]*scale + trans;
}
}
}
}
}
-#ifdef GPU
-cl_kernel get_crop_kernel()
-{
- static int init = 0;
- static cl_kernel kernel;
- if(!init){
- kernel = get_kernel("src/crop_layer.cl", "forward", 0);
- init = 1;
- }
- return kernel;
-}
-
-void forward_crop_layer_gpu(crop_layer layer, cl_mem input)
-{
- int flip = (layer.flip && rand()%2);
- int dh = rand()%(layer.h - layer.crop_height);
- int dw = rand()%(layer.w - layer.crop_width);
- int size = layer.batch*layer.c*layer.crop_width*layer.crop_height;
-
- cl_kernel kernel = get_crop_kernel();
- cl_command_queue queue = cl.queue;
-
- cl_uint i = 0;
- cl.error = clSetKernelArg(kernel, i++, sizeof(input), (void*) &input);
- cl.error = clSetKernelArg(kernel, i++, sizeof(layer.c), (void*) &layer.c);
- cl.error = clSetKernelArg(kernel, i++, sizeof(layer.h), (void*) &layer.h);
- cl.error = clSetKernelArg(kernel, i++, sizeof(layer.w), (void*) &layer.w);
- cl.error = clSetKernelArg(kernel, i++, sizeof(layer.crop_height), (void*) &layer.crop_height);
- cl.error = clSetKernelArg(kernel, i++, sizeof(layer.crop_width), (void*) &layer.crop_width);
- cl.error = clSetKernelArg(kernel, i++, sizeof(dh), (void*) &dh);
- cl.error = clSetKernelArg(kernel, i++, sizeof(dw), (void*) &dw);
- cl.error = clSetKernelArg(kernel, i++, sizeof(flip), (void*) &flip);
- cl.error = clSetKernelArg(kernel, i++, sizeof(layer.output_cl), (void*) &layer.output_cl);
- check_error(cl);
-
- const size_t global_size[] = {size};
-
- cl.error = clEnqueueNDRangeKernel(queue, kernel, 1, 0, global_size, 0, 0, 0, 0);
- check_error(cl);
-}
-
-#endif
--
Gitblit v1.10.0