diff options
Diffstat (limited to 'src/kernels/levelx/xconvgemm_part1.opencl')
| -rw-r--r-- | src/kernels/levelx/xconvgemm_part1.opencl | 33 |
1 files changed, 23 insertions, 10 deletions
diff --git a/src/kernels/levelx/xconvgemm_part1.opencl b/src/kernels/levelx/xconvgemm_part1.opencl index abdb5324..25ccba51 100644 --- a/src/kernels/levelx/xconvgemm_part1.opencl +++ b/src/kernels/levelx/xconvgemm_part1.opencl @@ -11,7 +11,6 @@ // uses parameters from the direct GEMM kernel. This is the part with the loads from memory (1/2). // This uses "CONVGEMM_WITH_IM2COL" as a switch to select between direct convgemm or first running // the im2col kernel to create a 'col' temporary matrix. -// TODO: Currently only works with 'CONVGEMM_WITH_IM2COL' set // // ================================================================================================= @@ -30,12 +29,17 @@ INLINE_FUNC real GlobalToPrivateCheckedImage(const __global real* restrict image 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 int dilation_h, const int dilation_w, + const bool kernel_flip) { // Im2col indices const int kernel_2d_index = kwg % (kernel_h * kernel_w); - const int kw_id = kernel_2d_index % kernel_w; - const int kh_id = kernel_2d_index / kernel_w; + const int kw_id = (kernel_flip) + ? kernel_w - kernel_2d_index % kernel_w - 1 + : kernel_2d_index % kernel_w; + const int kh_id = (kernel_flip) + ? kernel_h - kernel_2d_index / kernel_w - 1 + : kernel_2d_index / kernel_w; const int c_id = kwg / (kernel_h * kernel_w); 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; @@ -55,14 +59,15 @@ INLINE_FUNC real GlobalToPrivateCheckedImage(const __global real* restrict image // Loads global off-chip memory into local (shared) memory on-chip. This function is specific for // loading the image input tensor. This includes a bounds check. -INLINE_FUNC real GlobalToLocalCheckedImage(const __global realMD* restrict imagegm, LOCAL_PTR real* alm, +INLINE_FUNC real GlobalToLocalCheckedImage(const __global real* restrict imagegm, LOCAL_PTR real* alm, const int image_offset_batch, - const int h_id, const int w_id, const int kwg, + const int output_w, const int kwg, const int input_h, const int input_w, const int channels, 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 int dilation_h, const int dilation_w, + const bool kernel_flip) { #if MDIMCD == MDIMAD const int la0 = get_local_id(0); const int la1 = get_local_id(1); @@ -82,10 +87,17 @@ INLINE_FUNC real GlobalToLocalCheckedImage(const __global realMD* restrict image int idm = mg + GetGroupID0()*WGD; int idk = kg + kwg; + const int w_id = idm % output_w; + const int h_id = idm / output_w; + // Im2col indices const int kernel_2d_index = idk % (kernel_h * kernel_w); - const int kw_id = kernel_2d_index % kernel_w; - const int kh_id = kernel_2d_index / kernel_w; + const int kw_id = (kernel_flip) + ? kernel_w - kernel_2d_index % kernel_w - 1 + : kernel_2d_index % kernel_w; + const int kh_id = (kernel_flip) + ? kernel_h - kernel_2d_index / kernel_w - 1 + : kernel_2d_index / kernel_w; const int c_id = idk / (kernel_h * kernel_w); 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; @@ -104,7 +116,8 @@ INLINE_FUNC real GlobalToLocalCheckedImage(const __global realMD* restrict image } } -#endif +#endif // defined(ROUTINE_CONVGEMM) && !defined(CONVGEMM_WITH_IM2COL) + // ================================================================================================= // End of the C++11 raw string literal |