时间戳

首页 > 解决方案 > OpenCL 内核通过批处理实现 im2col

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

问题描述

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

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

我应该对 OpenCL 内核进行哪些更改以使其在 GPU 上的工作与在 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;
      }
}

标签： cimage-processingdeep-learningopencl

解决方案

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

您需要传入批量大小“B”，然后按批量大小将此副本更改为块副本（或只是循环）：

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;

推荐阅读