OpenCL Kernel implementing im2col with batch

Question

I am trying to adapt a secuential function writen for CPU to an OpenCL kernel for GPU. The function is the well known im2col used in many deep learning applications.

I have found some code on the OpenCV repository implementing this im2col function written in OpenCL but the one that I have to adapt uses a batch that confuses me and seems to be a bit different.

What should I change on the OpenCL kernel to make it work the same on GPU as it does on the CPU function?

CPU code

int fn_im2col_cpu(int I, int WI, int HI, int B, int KW, int KH, int WO, int HO, int PW, int PH, int SW, int SH, type *in_ptr, type *out_ptr) {
    PROFILING_HEADER_EXTERN(im2col);
    PROFILING_DEVICE(im2col, DEV_CPU);

    int i;                                             // scrolls input channels
    int w;                                             // scrolls channel columns (width)
    int h;                                             // scrolls channel rows (height)
    int kw;                                            // scrolls filter columns (width)
    int kh;                                            // scrolls filter rows (height)

 // we sweep all output pixels, and for each pixel we compute the associated input pixel
 #pragma omp parallel for private (kh, kw, h, w)
    for (i = 0; i < I; i++) {
        size_t out_addr = ((size_t)B * (size_t)WO * (size_t)HO * (size_t)KW * (size_t)KH * (size_t)i);
        size_t in_addr1 = (size_t)i * (size_t)B * (size_t)WI * (size_t)HI;
        for (kh = 0; kh < KH; kh++) {
            for (kw = 0; kw < KW; kw++) {
                for (h = 0; h < HO; h++) {
                    int hi = h * SH - PH + kh;
                    size_t in_addr2 = in_addr1 + ((size_t)hi * (size_t)B * (size_t)WI);
                    for (w = 0; w < WO; w++) {
                        int wi = w * SW - PW + kw;
                        int force_padding = (wi < 0) || (wi >= WI) || (hi < 0) || (hi >= HI);
                        if (force_padding) {
                            bzero(&out_ptr[out_addr], B*sizeof(type));
                        } else {
                            int in_addr = in_addr2 + (wi * B);
                            memcpy(&out_ptr[out_addr], &in_ptr[in_addr], B*sizeof(type));
                        }
                        out_addr+=B;
                    }
                }
            }
        }
    }
    return 1;
}

OpenCL kernel from https://github.com/opencv/opencv/blob/master/modules/dnn/src/opencl/im2col.cl

__kernel void im2col(__global const float *im_src, int im_src_offset,
                     int channels, int height_inp, int width_inp,
                     int kernel_h, int kernel_w, int pad_h, int pad_w, 
                     int stride_h, int stride_w,
                     int height_out, int width_out,
                     __global float *im_col, int im_col_offset
                    )
{
    int index = get_global_id(0);
    if (index >= height_out * width_out * channels)
        return;

    int j_out = index % width_out;
    int i_out = (index / width_out) % height_out;
    int c_inp = (index / width_out) / height_out;

    int c_out = c_inp * kernel_h * kernel_w;
    int i_inp = i_out * stride_h - pad_h;
    int j_inp = j_out * stride_w - pad_w;

    im_src += (c_inp * height_inp + i_inp) * width_inp + j_inp + im_src_offset;
    im_col += (c_out * height_out + i_out) * width_out + j_out + im_col_offset;

    for (int ki = 0; ki < kernel_h; ++ki)
        for (int kj = 0; kj < kernel_w; ++kj) {
            int i = i_inp + ki;
            int j = j_inp + kj;
            *im_col = (i >= 0 && j >= 0 && i < height_inp && j < width_inp) ?
                im_src[ki * width_inp + kj] : 0;
            im_col += height_out * width_out;
      }
}

What do you mean by "the one that I have to adapt uses a batch that confuses me"? — pmdj

Hashman Hashman · Accepted Answer · 2019-12-07T06:11:30

Your C version folds the batch into the lowest dimension. The opencl version isn't even using batch.

You need to pass in the batch size "B", and change this copy to a block copy (or just do a loop over) by the batch size:

for (int b=0; b<B; b++) *(im_col*B+b) = (i >= 0 && j >= 0 && i < height_inp && j < width_inp) ? im_src[(ki * width_inp + kj)*B + b] : 0;

to emulate the memcpy(..., B*sizeof(type)).

And then just stride B times more:

im_col += height_out * width_out * B;

OpenCL Kernel implementing im2col with batch

1 Answers