diff options
author | Cedric Nugteren <web@cedricnugteren.nl> | 2017-08-24 21:11:12 +0200 |
---|---|---|
committer | Cedric Nugteren <web@cedricnugteren.nl> | 2017-08-24 21:11:12 +0200 |
commit | 4d9d03ba512fb404537180a42cf17fb3e10a8033 (patch) | |
tree | 8205f3939bdaea2c22af48b8adc0d83a85dd937b /src | |
parent | a8c26594d9ee5b735b9d1b0400ac7eba4abf02fb (diff) |
Completed im2col implementation
Diffstat (limited to 'src')
-rw-r--r-- | src/kernels/levelx/im2col.opencl | 17 | ||||
-rw-r--r-- | src/routines/levelx/xim2col.cpp | 10 |
2 files changed, 16 insertions, 11 deletions
diff --git a/src/kernels/levelx/im2col.opencl b/src/kernels/levelx/im2col.opencl index 3f10881f..a141db41 100644 --- a/src/kernels/levelx/im2col.opencl +++ b/src/kernels/levelx/im2col.opencl @@ -15,12 +15,17 @@ // literal). Comment-out this line for syntax-highlighting when developing. R"( -// ================================================================================================= +// Work-group size parameters re-used from the 'copy' kernel +#ifndef COPY_DIMX + #define COPY_DIMX 8 // Local workgroup size in the first dimension (x) +#endif +#ifndef COPY_DIMY + #define COPY_DIMY 8 // Local workgroup size in the second dimension (y) +#endif -#define WGS1 16 -#define WGS2 16 +// ================================================================================================= -__kernel __attribute__((reqd_work_group_size(WGS1, WGS2, 1))) +__kernel __attribute__((reqd_work_group_size(COPY_DIMX, COPY_DIMY, 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, @@ -31,8 +36,8 @@ void im2col(const int input_h, const int input_w, __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 w_id = get_global_id(0); // image width, max 'output_w' + const int h_id = get_global_id(1); // image height, max 'output_h' const int c_id = get_global_id(2); // input channels if (h_id < output_h && w_id < output_w) { diff --git a/src/routines/levelx/xim2col.cpp b/src/routines/levelx/xim2col.cpp index 10c9c10c..51171eb5 100644 --- a/src/routines/levelx/xim2col.cpp +++ b/src/routines/levelx/xim2col.cpp @@ -22,7 +22,7 @@ namespace clblast { // Constructor: forwards to base class constructor template <typename T> Xim2col<T>::Xim2col(Queue &queue, EventPointer event, const std::string &name): - Routine(queue, event, name, {}, PrecisionValue<T>(), {}, { + Routine(queue, event, name, {"Copy"}, PrecisionValue<T>(), {}, { #include "../../kernels/levelx/im2col.opencl" }) { } @@ -71,10 +71,10 @@ void Xim2col<T>::DoIm2col(const size_t channels, const size_t height, const size kernel.SetArgument(15, static_cast<int>(col_offset)); // Launches the kernel - const auto h_ceiled = Ceil(output_h, 16); - const auto w_ceiled = Ceil(output_w, 16); - auto global = std::vector<size_t>{h_ceiled, w_ceiled, channels}; - auto local = std::vector<size_t>{16, 16, 1}; + const auto w_ceiled = Ceil(output_w, db_["COPY_DIMY"]); + const auto h_ceiled = Ceil(output_h, db_["COPY_DIMX"]); + const auto global = std::vector<size_t>{w_ceiled, h_ceiled, channels}; + const auto local = std::vector<size_t>{db_["COPY_DIMX"], db_["COPY_DIMY"], 1}; RunKernel(kernel, queue_, device_, global, local, event_); } |