c - OpenCL 内核通过批处理实现 im2col

Question

我正在尝试将为 CPU 编写的连续函数改编为 GPU 的 OpenCL 内核。该函数是许多深度学习应用中使用的众所周知的 im2col。

我在 OpenCV 存储库上找到了一些代码，实现了用 OpenCL 编写的 im2col 函数，但我必须调整的代码使用了一批让我感到困惑的代码，并且似乎有点不同。

我应该对 OpenCL 内核进行哪些更改，以使其在 GPU 上的工作方式与在 CPU 功能上的工作方式相同？

CPU代码

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;
}

来自 https://github.com/opencv/opencv/blob/master/modules/dnn/src/opencl/im2col.cl 的 OpenCL 内核

__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;
      }
}

Answer 1

您的 C 版本将批处理折叠到最低维度。 opencl 版本甚至没有使用批处理。

您需要传入批量大小“B”，并按批量大小将此副本更改为 block 副本(或只是进行循环):

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;

模拟 memcpy(..., B*sizeof(type))。

然后再跨步 B 倍:

im_col += height_out * width_out * B;