diff options
author | Cedric Nugteren <web@cedricnugteren.nl> | 2018-05-19 21:02:44 +0200 |
---|---|---|
committer | Cedric Nugteren <web@cedricnugteren.nl> | 2018-05-19 21:02:44 +0200 |
commit | 27b52ac2c886238de59253a065c1fbbb970c73f1 (patch) | |
tree | 32f93133eff6e76be45e1e0ca031b4afb09519be /src/kernels/levelx/xconvgemm_part1.opencl | |
parent | cbcd4ff7e8e21584a9a1f405c9f4cb979a73b718 (diff) |
Second version of direct reading from image tensor for convgemm: also with local memory support now
Diffstat (limited to 'src/kernels/levelx/xconvgemm_part1.opencl')
-rw-r--r-- | src/kernels/levelx/xconvgemm_part1.opencl | 110 |
1 files changed, 110 insertions, 0 deletions
diff --git a/src/kernels/levelx/xconvgemm_part1.opencl b/src/kernels/levelx/xconvgemm_part1.opencl new file mode 100644 index 00000000..ac13219f --- /dev/null +++ b/src/kernels/levelx/xconvgemm_part1.opencl @@ -0,0 +1,110 @@ + +// ================================================================================================= +// 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 an implementation of 3D convolution on a 4D image using GEMM kernels. It +// uses parameters from the direct GEMM kernel. This is the part with the loads from memory (1/2). +// +// ================================================================================================= + +// 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"( + +// ================================================================================================= +#if defined(ROUTINE_CONVGEMM) + +// Loads global off-chip memory into thread-private register files. This function is specific for +// loading the image input tensor. This includes a bounds check. +INLINE_FUNC real GlobalToPrivateCheckedImage(const __global real* restrict imagegm, const int image_offset_batch, + const int h_id, const int w_id, 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) { + + // 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 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; + + // With bounds check + real result; + if (h_index >= 0 && h_index < input_h && + w_index >= 0 && w_index < input_w) { + const int image_index = w_index + input_w * (h_index + input_h * c_id); + result = imagegm[image_index + image_offset_batch]; + } + else { + SetToZero(result); + } + return result; +} + +// 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, + const int image_offset_batch, + const int h_id, const int w_id, 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) { + #if MDIMCD == MDIMAD + const int la0 = get_local_id(0); + const int la1 = get_local_id(1); + #else + const int tid = get_local_id(0) + MDIMCD*get_local_id(1); + const int la0 = tid % MDIMAD; + const int la1 = tid / MDIMAD; + #endif + #pragma unroll + for (int _mia = 0; _mia < MWAD; _mia += 1) { + #pragma unroll + for (int _kia = 0; _kia < KWAD; _kia += 1) { + + // Computes the indices for the global memory + int mg = _mia + la0*MWAD; + int kg = _kia + la1*KWAD; + int idm = mg + GetGroupID0()*WGD; + int idk = kg + kwg; + + // 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 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; + + // Loads the data from global memory into the local memory + if (h_index >= 0 && h_index < input_h && + w_index >= 0 && w_index < input_w) { + const int image_index = w_index + input_w * (h_index + input_h * c_id); + const real result = imagegm[image_index + image_offset_batch]; + alm[kg*(WGD + PADA) + mg] = result; + } + else { + SetToZero(alm[kg*(WGD + PADA) + mg]); + } + } + } +} + +#endif +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= |