summaryrefslogtreecommitdiff
path: root/src/kernels/levelx/im2col.opencl
blob: 3f10881f472731c113516b58c8e5ec7675837f14 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
// =================================================================================================
// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This
// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max-
// width of 100 characters per line.
//
// Author(s):
//   Cedric Nugteren <www.cedricnugteren.nl>
//
// This file contains the im2col kernel.
//
// =================================================================================================

// Enables loading of this file using the C++ pre-processor's #include (C++11 standard raw string
// literal). Comment-out this line for syntax-highlighting when developing.
R"(

// =================================================================================================

#define WGS1 16
#define WGS2 16

__kernel __attribute__((reqd_work_group_size(WGS1, WGS2, 1)))
void im2col(const int input_h, const int input_w,
            const int output_h, const int output_w,
            const int kernel_h, const int kernel_w,
            const int pad_h, const int pad_w,
            const int stride_h, const int stride_w,
            const int dilation_h, const int dilation_w,
            const __global real* restrict im_buffer, const int im_offset,
            __global real* col_buffer, const int col_offset) {

  // Thread IDs
  const int h_id = get_global_id(0); // image height, max 'output_h'
  const int w_id = get_global_id(1); // image width, max 'output_w'
  const int c_id = get_global_id(2); // input channels
  if (h_id < output_h && w_id < output_w) {

    #pragma unroll
    for (int kh_id = 0; kh_id < kernel_h; ++kh_id) { // kernel height
      #pragma unroll
      for (int kw_id = 0; kw_id < kernel_w; ++kw_id) { // kernel width

        // Retrieves the input value
        const int h_index = -pad_h + kh_id * dilation_h + stride_h * h_id;
        const int w_index = -pad_w + kw_id * dilation_w + stride_w * w_id;
        real val;
        if (h_index >= 0 && h_index < input_h &&
            w_index >= 0 && w_index < input_w) {
          const int input_index = w_index + input_w * (h_index + input_h * c_id);
          val = im_buffer[input_index + im_offset];
        }
        else {
          SetToZero(val);
        }

        // Sets the output value
        const int kernel_index = kw_id + kernel_w * kh_id;
        const int patch_index = w_id + output_w * h_id;
        const int output_index = patch_index + kernel_index * output_w * output_h +
                                  c_id * output_w * output_h * kernel_h * kernel_w;
        col_buffer[output_index + col_offset] = val;
      }
    }
  }
}

// =================================================================================================

// End of the C++11 raw string literal
)"

// =================================================================================================