Second version of direct reading from image tensor for convgemm: also with local memory support now

3 files changed, 122 insertions, 33 deletions

diff --git a/src/kernels/level3/xconvgemm.opencl b/src/kernels/level3/xconvgemm.opencl
index cddb6785..7824f958 100644
--- a/src/kernels/level3/xconvgemm.opencl
+++ b/src/kernels/level3/xconvgemm.opencl
@@ -19,35 +19,6 @@ 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) {
-  real result;
-
-  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;
-  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;
-}
-
 // ConvGEMM kernel
 __kernel __attribute__((reqd_work_group_size(MDIMCD, NDIMCD, 1)))
 void Xconvgemm(const int num_patches, const int num_kernels, const int patch_size,
@@ -189,8 +160,15 @@ void Xconvgemm(const int num_patches, const int num_kernels, const int patch_siz
     int kwg = 0;
     for (; kwg < (patch_size/WGD) * WGD; kwg+=WGD) {
 
-      // Loads data: off-chip --> local (matrix A and B)
-      GlobalToLocalCheckedA(colgms, alm, num_patches, col_offset_batch, kwg, false, false, num_patches, patch_size);
+      // Loads data: off-chip --> local
+      #if defined(CONVGEMM_WITH_IM2COL)
+        GlobalToLocalCheckedA(colgms, alm, num_patches, col_offset_batch, kwg, false, false, num_patches, patch_size);
+      #else
+        GlobalToLocalCheckedImage(imagegm, alm, image_offset_batch, h_id, w_id, kwg,
+                                  input_h, input_w, channels, kernel_h, kernel_w,
+                                  pad_h, pad_w, stride_h, stride_w,
+                                  dilation_h, dilation_w);
+      #endif
       GlobalToLocalCheckedB(kernelgms, blm, patch_size, kernel_offset, kwg, true, false, num_kernels, patch_size);
       barrier(CLK_LOCAL_MEM_FENCE);
 
@@ -200,7 +178,7 @@ void Xconvgemm(const int num_patches, const int num_kernels, const int patch_siz
         for (int _pit = 0; _pit < KWID; _pit += 1) {
           int kg = pwi + _pit;
 
-          // Loads data: local --> private (matrix A and B)
+          // Loads data: local --> private
           #pragma unroll
           for (int _mi = 0; _mi < MWID; _mi += 1) {
             apd[_mi] = LocalToPrivateDirectA(alm, _mi, kg, false);
@@ -226,7 +204,7 @@ void Xconvgemm(const int num_patches, const int num_kernels, const int patch_siz
     // Loop over the remaining part (incomplete tile in K-dimension)
     for (; kwg < patch_size; ++kwg) {
 
-      // Loads data: off-chip --> private (matrix A and B)
+      // Loads data: off-chip --> private
       #pragma unroll
       for (int _mi = 0; _mi < MWID; _mi += 1) {
         apd[_mi] = GlobalToPrivateCheckedImage(imagegm, image_offset_batch, h_id, w_id, kwg,
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
+)"
+
+// =================================================================================================
diff --git a/src/routines/levelx/xconvgemm.cpp b/src/routines/levelx/xconvgemm.cpp
index 23335261..b1f07809 100644
--- a/src/routines/levelx/xconvgemm.cpp
+++ b/src/routines/levelx/xconvgemm.cpp
@@ -31,6 +31,7 @@ Xconvgemm<T>::Xconvgemm(Queue &queue, EventPointer event, const std::string &nam
             #include "../../kernels/level3/xgemm_direct_part2.opencl"
             #include "../../kernels/level3/xgemm_direct_part3.opencl"
             , // separated in multiple parts to prevent C1091 in MSVC 2013
+            #include "../../kernels/levelx/xconvgemm_part1.opencl"
             #include "../../kernels/level3/xconvgemm.opencl"
         }) {
 }