Code cleaning & training new YOLO model

Edmond Yoo

2018-09-17 504ece5b00f192d5c1b343fd06ce1648f9139180

 src/col2im_kernels.cu

@@ -1,3 +1,7 @@
#include "cuda_runtime.h"
#include "curand.h"
#include "cublas_v2.h"

extern "C" {
#include "col2im.h"
#include "cuda.h"
@@ -33,16 +37,15 @@
                val += data_col[offset + h_col * coeff_h_col + w_col * coeff_w_col];
            }
        }
        data_im[index] = val;
        data_im[index] += val;
    }
}

void col2im_ongpu(float *im,
void col2im_ongpu(float *data_col,
        int channels, int height, int width,
        int ksize, int stride, int pad, float *data_col){
        int ksize, int stride, int pad, float *data_im){
    // We are going to launch channels * height_col * width_col kernels, each
    // kernel responsible for copying a single-channel grid.
    pad = pad ? ksize/2 : 0;
    int height_col = (height + 2 * pad - ksize) / stride + 1;
    int width_col = (width + 2 * pad - ksize) / stride + 1;
    int num_kernels = channels * height * width;
@@ -50,65 +53,6 @@
        BLOCK>>>(
                num_kernels, data_col, height, width, ksize, pad,
                stride, height_col,
                width_col, im);
                width_col, data_im);
}

/*
   __global__ void col2im_kernel(float *data_col,
   int channels, int height, int width,
   int ksize, int stride, int pad, float *data_im)
   {

   int height_col = (height - ksize) / stride + 1;
   int width_col = (width - ksize) / stride + 1;
   if (pad){
   height_col = 1 + (height-1) / stride;
   width_col = 1 + (width-1) / stride;
   pad = ksize/2;
   }

   int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
   if(id >= channels*height*width) return;

   int index = id;
   int w = id%width + pad;
   id /= width;
   int h = id%height + pad;
   id /= height;
   int c = id%channels;

   int w_start = (w-ksize+stride)/stride;
   int w_end = w/stride + 1;

   int h_start = (h-ksize+stride)/stride;
   int h_end = h/stride + 1;

// int rows = channels * ksize * ksize;
// int cols = height_col*width_col;
int col_offset = (c*ksize*ksize + h * ksize + w)*height_col*width_col;
int h_coeff = (1-stride*ksize*height_col)*width_col;
int w_coeff = 1-stride*height_col*width_col;
float val = 0;
int h_col, w_col;
for(h_col = h_start; h_col < h_end; ++h_col){
for(w_col = w_start; w_col < w_end; ++w_col){
int col_index = col_offset +h_col*h_coeff + w_col*w_coeff;
float part = (w_col < 0 || h_col < 0 || h_col >= height_col || w_col >= width_col) ? 0 : data_col[col_index];
val += part;
}
}
data_im[index] = val;
}


extern "C" void col2im_ongpu(float *data_col,
int channels,  int height,  int width,
int ksize,  int stride,  int pad, float *data_im)
{

size_t n = channels*height*width;

col2im_kernel<<<cuda_gridsize(n), BLOCK>>>(data_col, channels, height, width, ksize, stride, pad, data_im);
check_error(cudaPeekAtLastError());
}
 */