Added xDOT/xDOTU/xDOTC dot-product routines

9 files changed, 618 insertions, 0 deletions

diff --git a/src/clblast.cc b/src/clblast.cc
index c99ad7b1..3303085e 100644
--- a/src/clblast.cc
+++ b/src/clblast.cc
@@ -22,6 +22,9 @@
 #include "internal/routines/level1/xscal.h"
 #include "internal/routines/level1/xcopy.h"
 #include "internal/routines/level1/xaxpy.h"
+#include "internal/routines/level1/xdot.h"
+#include "internal/routines/level1/xdotu.h"
+#include "internal/routines/level1/xdotc.h"
 
 // BLAS level-2 includes
 #include "internal/routines/level2/xgemv.h"
@@ -177,6 +180,90 @@ template StatusCode Axpy<double2>(const size_t,
                                   cl_mem, const size_t, const size_t,
                                   cl_command_queue* queue, cl_event* event);
 
+// DOT
+template <typename T>
+StatusCode Dot(const size_t n,
+               cl_mem dot_buffer, const size_t dot_offset,
+               const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+               const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+               cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xdot<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoDot(n,
+                       Buffer<T>(dot_buffer), dot_offset,
+                       Buffer<T>(x_buffer), x_offset, x_inc,
+                       Buffer<T>(y_buffer), y_offset, y_inc);
+}
+template StatusCode Dot<float>(const size_t,
+                               cl_mem, const size_t,
+                               const cl_mem, const size_t, const size_t,
+                               const cl_mem, const size_t, const size_t,
+                               cl_command_queue* queue, cl_event* event);
+template StatusCode Dot<double>(const size_t,
+                                cl_mem, const size_t,
+                                const cl_mem, const size_t, const size_t,
+                                const cl_mem, const size_t, const size_t,
+                                cl_command_queue* queue, cl_event* event);
+
+// DOTU
+template <typename T>
+StatusCode Dotu(const size_t n,
+                cl_mem dot_buffer, const size_t dot_offset,
+                const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+                cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xdotu<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoDotu(n,
+                        Buffer<T>(dot_buffer), dot_offset,
+                        Buffer<T>(x_buffer), x_offset, x_inc,
+                        Buffer<T>(y_buffer), y_offset, y_inc);
+}
+template StatusCode Dotu<float2>(const size_t,
+                                 cl_mem, const size_t,
+                                 const cl_mem, const size_t, const size_t,
+                                 const cl_mem, const size_t, const size_t,
+                                 cl_command_queue* queue, cl_event* event);
+template StatusCode Dotu<double2>(const size_t,
+                                  cl_mem, const size_t,
+                                  const cl_mem, const size_t, const size_t,
+                                  const cl_mem, const size_t, const size_t,
+                                  cl_command_queue* queue, cl_event* event);
+
+// DOTC
+template <typename T>
+StatusCode Dotc(const size_t n,
+                cl_mem dot_buffer, const size_t dot_offset,
+                const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+                cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xdotc<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoDotc(n,
+                        Buffer<T>(dot_buffer), dot_offset,
+                        Buffer<T>(x_buffer), x_offset, x_inc,
+                        Buffer<T>(y_buffer), y_offset, y_inc);
+}
+template StatusCode Dotc<float2>(const size_t,
+                                 cl_mem, const size_t,
+                                 const cl_mem, const size_t, const size_t,
+                                 const cl_mem, const size_t, const size_t,
+                                 cl_command_queue* queue, cl_event* event);
+template StatusCode Dotc<double2>(const size_t,
+                                  cl_mem, const size_t,
+                                  const cl_mem, const size_t, const size_t,
+                                  const cl_mem, const size_t, const size_t,
+                                  cl_command_queue* queue, cl_event* event);
+
 // =================================================================================================
 // BLAS level-2 (matrix-vector) routines
 // =================================================================================================
diff --git a/src/clblast_c.cc b/src/clblast_c.cc
index ab3994fb..eccf517f 100644
--- a/src/clblast_c.cc
+++ b/src/clblast_c.cc
@@ -203,6 +203,84 @@ StatusCode CLBlastZaxpy(const size_t n,
   return static_cast<StatusCode>(status);
 }
 
+// DOT
+StatusCode CLBlastSdot(const size_t n,
+                       cl_mem dot_buffer, const size_t dot_offset,
+                       const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                       const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+                       cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Dot<float>(n,
+                                    dot_buffer, dot_offset,
+                                    x_buffer, x_offset, x_inc,
+                                    y_buffer, y_offset, y_inc,
+                                    queue, event);
+  return static_cast<StatusCode>(status);
+}
+StatusCode CLBlastDdot(const size_t n,
+                       cl_mem dot_buffer, const size_t dot_offset,
+                       const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                       const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+                       cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Dot<double>(n,
+                                     dot_buffer, dot_offset,
+                                     x_buffer, x_offset, x_inc,
+                                     y_buffer, y_offset, y_inc,
+                                     queue, event);
+  return static_cast<StatusCode>(status);
+}
+
+// DOTU
+StatusCode CLBlastCdotu(const size_t n,
+                        cl_mem dot_buffer, const size_t dot_offset,
+                        const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                        const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+                        cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Dotu<float2>(n,
+                                      dot_buffer, dot_offset,
+                                      x_buffer, x_offset, x_inc,
+                                      y_buffer, y_offset, y_inc,
+                                      queue, event);
+  return static_cast<StatusCode>(status);
+}
+StatusCode CLBlastZdotu(const size_t n,
+                        cl_mem dot_buffer, const size_t dot_offset,
+                        const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                        const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+                        cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Dotu<double2>(n,
+                                       dot_buffer, dot_offset,
+                                       x_buffer, x_offset, x_inc,
+                                       y_buffer, y_offset, y_inc,
+                                       queue, event);
+  return static_cast<StatusCode>(status);
+}
+
+// DOTC
+StatusCode CLBlastCdotc(const size_t n,
+                        cl_mem dot_buffer, const size_t dot_offset,
+                        const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                        const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+                        cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Dotc<float2>(n,
+                                      dot_buffer, dot_offset,
+                                      x_buffer, x_offset, x_inc,
+                                      y_buffer, y_offset, y_inc,
+                                      queue, event);
+  return static_cast<StatusCode>(status);
+}
+StatusCode CLBlastZdotc(const size_t n,
+                        cl_mem dot_buffer, const size_t dot_offset,
+                        const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                        const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
+                        cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Dotc<double2>(n,
+                                       dot_buffer, dot_offset,
+                                       x_buffer, x_offset, x_inc,
+                                       y_buffer, y_offset, y_inc,
+                                       queue, event);
+  return static_cast<StatusCode>(status);
+}
+
 // =================================================================================================
 // BLAS level-2 (matrix-vector) routines
 // =================================================================================================
diff --git a/src/database.cc b/src/database.cc
index 258d861e..b7275dad 100644
--- a/src/database.cc
+++ b/src/database.cc
@@ -13,6 +13,7 @@
 
 #include "internal/database.h"
 #include "internal/database/xaxpy.h"
+#include "internal/database/xdot.h"
 #include "internal/database/xgemv.h"
 #include "internal/database/xgemm.h"
 #include "internal/database/copy.h"
@@ -28,6 +29,7 @@ namespace clblast {
 // Initializes the database
 const std::vector<Database::DatabaseEntry> Database::database = {
   XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble,
+  XdotSingle, XdotDouble, XdotComplexSingle, XdotComplexDouble,
   XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble,
   XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble,
   CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble,
diff --git a/src/kernels/common.opencl b/src/kernels/common.opencl
index 8e71429e..f2a2e7a7 100644
--- a/src/kernels/common.opencl
+++ b/src/kernels/common.opencl
@@ -109,6 +109,13 @@ R"(
   #define SetToOne(a) a = ONE
 #endif
 
+// Adds two complex variables
+#if PRECISION == 3232 || PRECISION == 6464
+  #define Add(c, a, b) c.x = a.x + b.x; c.y = a.y + b.y
+#else
+  #define Add(c, a, b) c = a + b
+#endif
+
 // Multiply two complex variables (used in the defines below)
 #if PRECISION == 3232 || PRECISION == 6464
   #define MulReal(a, b) a.x*b.x - a.y*b.y
diff --git a/src/kernels/level1/xdot.opencl b/src/kernels/level1/xdot.opencl
new file mode 100644
index 00000000..e13eb3c1
--- /dev/null
+++ b/src/kernels/level1/xdot.opencl
@@ -0,0 +1,106 @@
+
+// =================================================================================================
+// 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 Xdot kernel. It implements a dot-product computation using reduction
+// kernels. Reduction is split in two parts. In the first (main) kernel the X and Y vectors are
+// multiplied, followed by a per-thread and a per-workgroup reduction. The second (epilogue) kernel
+// is executed with a single workgroup only, computing the final result.
+//
+// =================================================================================================
+
+// 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"(
+
+// Parameters set by the tuner or by the database. Here they are given a basic default value in case
+// this kernel file is used outside of the CLBlast library.
+#ifndef WGS1
+  #define WGS1 64     // The local work-group size of the main kernel
+#endif
+#ifndef WGS2
+  #define WGS2 64     // The local work-group size of the epilogue kernel
+#endif
+
+// =================================================================================================
+
+// The main reduction kernel, performing the multiplication and the majority of the sum operation
+__attribute__((reqd_work_group_size(WGS1, 1, 1)))
+__kernel void Xdot(const int n,
+                   const __global real* restrict xgm, const int x_offset, const int x_inc,
+                   const __global real* restrict ygm, const int y_offset, const int y_inc,
+                   __global real* output, const int do_conjugate) {
+  __local real lm[WGS1];
+  const int lid = get_local_id(0);
+  const int wgid = get_group_id(0);
+  const int num_groups = get_num_groups(0);
+
+  // Performs multiplication and the first steps of the reduction
+  real acc;
+  SetToZero(acc);
+  int id = wgid*WGS1 + lid;
+  while (id < n) {
+    real x = xgm[id*x_inc + x_offset];
+    real y = ygm[id*y_inc + y_offset];
+    if (do_conjugate) { COMPLEX_CONJUGATE(x); }
+    MultiplyAdd(acc, x, y);
+    id += WGS1*num_groups;
+  }
+  lm[lid] = acc;
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  // Performs reduction in local memory
+  #pragma unroll
+  for (int s=WGS1/2; s>0; s=s>>1) {
+    if (lid < s) {
+      Add(lm[lid], lm[lid], lm[lid + s]);
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+  }
+
+  // Stores the per-workgroup result
+  if (lid == 0) {
+    output[wgid] = lm[0];
+  }
+}
+
+// =================================================================================================
+
+// The epilogue reduction kernel, performing the final bit of the sum operation. This kernel has to
+// be launched with a single workgroup only.
+__attribute__((reqd_work_group_size(WGS2, 1, 1)))
+__kernel void XdotEpilogue(const __global real* restrict input,
+                           __global real* dot, const int dot_offset) {
+  __local real lm[WGS2];
+  const int lid = get_local_id(0);
+
+  // Performs the first step of the reduction while loading the data
+  Add(lm[lid], input[lid], input[lid + WGS2]);
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  // Performs reduction in local memory
+  #pragma unroll
+  for (int s=WGS2/2; s>0; s=s>>1) {
+    if (lid < s) {
+      Add(lm[lid], lm[lid], lm[lid + s]);
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+  }
+
+  // Stores the final result
+  if (lid == 0) {
+    dot[dot_offset] = lm[0];
+  }
+}
+
+// =================================================================================================
+
+// End of the C++11 raw string literal
+)"
+
+// =================================================================================================
diff --git a/src/routines/level1/xdot.cc b/src/routines/level1/xdot.cc
new file mode 100644
index 00000000..a0c1e756
--- /dev/null
+++ b/src/routines/level1/xdot.cc
@@ -0,0 +1,115 @@
+
+// =================================================================================================
+// 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 implements the Xdot class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level1/xdot.h"
+
+#include <string>
+#include <vector>
+
+namespace clblast {
+// =================================================================================================
+
+// Specific implementations to get the memory-type based on a template argument
+template <> const Precision Xdot<float>::precision_ = Precision::kSingle;
+template <> const Precision Xdot<double>::precision_ = Precision::kDouble;
+template <> const Precision Xdot<float2>::precision_ = Precision::kComplexSingle;
+template <> const Precision Xdot<double2>::precision_ = Precision::kComplexDouble;
+
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xdot<T>::Xdot(Queue &queue, Event &event, const std::string &name):
+    Routine<T>(queue, event, name, {"Xdot"}, precision_) {
+  source_string_ =
+    #include "../../kernels/level1/xdot.opencl"
+  ;
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xdot<T>::DoDot(const size_t n,
+                          const Buffer<T> &dot_buffer, const size_t dot_offset,
+                          const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                          const Buffer<T> &y_buffer, const size_t y_offset, const size_t y_inc,
+                          const bool do_conjugate) {
+
+  // Makes sure all dimensions are larger than zero
+  if (n == 0) { return StatusCode::kInvalidDimension; }
+
+  // Tests the vectors for validity
+  auto status = TestVectorX(n, x_buffer, x_offset, x_inc, sizeof(T));
+  if (ErrorIn(status)) { return status; }
+  status = TestVectorY(n, y_buffer, y_offset, y_inc, sizeof(T));
+  if (ErrorIn(status)) { return status; }
+  status = TestVectorDot(1, dot_buffer, dot_offset, sizeof(T));
+  if (ErrorIn(status)) { return status; }
+
+  // Retrieves the Xdot kernels from the compiled binary
+  try {
+    auto& program = GetProgramFromCache();
+    auto kernel1 = Kernel(program, "Xdot");
+    auto kernel2 = Kernel(program, "XdotEpilogue");
+
+    // Creates the buffer for intermediate values
+    auto temp_size = 2*db_["WGS2"];
+    auto temp_buffer = Buffer<T>(context_, temp_size);
+
+    // Sets the kernel arguments
+    kernel1.SetArgument(0, static_cast<int>(n));
+    kernel1.SetArgument(1, x_buffer());
+    kernel1.SetArgument(2, static_cast<int>(x_offset));
+    kernel1.SetArgument(3, static_cast<int>(x_inc));
+    kernel1.SetArgument(4, y_buffer());
+    kernel1.SetArgument(5, static_cast<int>(y_offset));
+    kernel1.SetArgument(6, static_cast<int>(y_inc));
+    kernel1.SetArgument(7, temp_buffer());
+    kernel1.SetArgument(8, static_cast<int>(do_conjugate));
+
+    // Launches the main kernel
+    auto global1 = std::vector<size_t>{db_["WGS1"]*temp_size};
+    auto local1 = std::vector<size_t>{db_["WGS1"]};
+    status = RunKernel(kernel1, global1, local1);
+    if (ErrorIn(status)) { return status; }
+
+    // Sets the arguments for the epilogue kernel
+    kernel2.SetArgument(0, temp_buffer());
+    kernel2.SetArgument(1, dot_buffer());
+    kernel2.SetArgument(2, static_cast<int>(dot_offset));
+
+    // Launches the epilogue kernel
+    auto global2 = std::vector<size_t>{db_["WGS2"]};
+    auto local2 = std::vector<size_t>{db_["WGS2"]};
+    status = RunKernel(kernel2, global2, local2);
+    if (ErrorIn(status)) { return status; }
+
+    // Waits for all kernels to finish
+    queue_.Finish();
+
+    // Succesfully finished the computation
+    return StatusCode::kSuccess;
+  } catch (...) { return StatusCode::kInvalidKernel; }
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xdot<float>;
+template class Xdot<double>;
+template class Xdot<float2>;
+template class Xdot<double2>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/routines/level1/xdotc.cc b/src/routines/level1/xdotc.cc
new file mode 100644
index 00000000..f414f556
--- /dev/null
+++ b/src/routines/level1/xdotc.cc
@@ -0,0 +1,49 @@
+
+// =================================================================================================
+// 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 implements the Xdotc class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level1/xdotc.h"
+
+#include <string>
+#include <vector>
+
+namespace clblast {
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xdotc<T>::Xdotc(Queue &queue, Event &event, const std::string &name):
+    Xdot<T>(queue, event, name) {
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xdotc<T>::DoDotc(const size_t n,
+                            const Buffer<T> &dot_buffer, const size_t dot_offset,
+                            const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                            const Buffer<T> &y_buffer, const size_t y_offset, const size_t y_inc) {
+  return DoDot(n, dot_buffer, dot_offset,
+               x_buffer, x_offset, x_inc,
+               y_buffer, y_offset, y_inc,
+               true);
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xdotc<float2>;
+template class Xdotc<double2>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/routines/level1/xdotu.cc b/src/routines/level1/xdotu.cc
new file mode 100644
index 00000000..0b1bd2a8
--- /dev/null
+++ b/src/routines/level1/xdotu.cc
@@ -0,0 +1,49 @@
+
+// =================================================================================================
+// 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 implements the Xdotu class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level1/xdotu.h"
+
+#include <string>
+#include <vector>
+
+namespace clblast {
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xdotu<T>::Xdotu(Queue &queue, Event &event, const std::string &name):
+    Xdot<T>(queue, event, name) {
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xdotu<T>::DoDotu(const size_t n,
+                            const Buffer<T> &dot_buffer, const size_t dot_offset,
+                            const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                            const Buffer<T> &y_buffer, const size_t y_offset, const size_t y_inc) {
+  return DoDot(n, dot_buffer, dot_offset,
+               x_buffer, x_offset, x_inc,
+               y_buffer, y_offset, y_inc,
+               false);
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xdotu<float2>;
+template class Xdotu<double2>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/tuning/xdot.cc b/src/tuning/xdot.cc
new file mode 100644
index 00000000..ff6bee16
--- /dev/null
+++ b/src/tuning/xdot.cc
@@ -0,0 +1,125 @@
+
+// =================================================================================================
+// 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 uses the CLTune auto-tuner to tune the xdot OpenCL kernels. Note that the results are
+// not verified, since the result is not final and depends on the WGS2 parameter.
+//
+// =================================================================================================
+
+#include <string>
+#include <vector>
+
+#include "internal/utilities.h"
+#include "internal/tuning.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class TuneXdot {
+ public:
+
+  // The representative kernel and the source code
+  static std::string KernelFamily() { return "xdot"; }
+  static std::string KernelName() { return "Xdot"; }
+  static std::string GetSources() {
+    return
+      #include "../src/kernels/common.opencl"
+      #include "../src/kernels/level1/xdot.opencl"
+    ;
+  }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() { return {kArgN}; }
+
+  // Tests for valid arguments
+  static void TestValidArguments(const Arguments<T> &) { }
+
+  // Sets the default values for the arguments
+  static size_t DefaultM() { return 1; } // N/A for this kernel
+  static size_t DefaultN() { return 4096*1024; }
+  static size_t DefaultK() { return 1; } // N/A for this kernel
+  static double DefaultFraction() { return 1.0; } // N/A for this kernel
+
+  // Describes how to obtain the sizes of the buffers
+  static size_t GetSizeX(const Arguments<T> &args) { return args.n; }
+  static size_t GetSizeY(const Arguments<T> &args) { return args.n; }
+  static size_t GetSizeA(const Arguments<T> &) { return 1; } // N/A for this kernel
+  static size_t GetSizeB(const Arguments<T> &) { return 1; } // N/A for this kernel
+  static size_t GetSizeC(const Arguments<T> &) { return 1; } // N/A for this kernel
+  static size_t GetSizeTemp(const Arguments<T> &args) { return args.n; } // Worst case
+
+  // Sets the tuning parameters and their possible values
+  static void SetParameters(cltune::Tuner &tuner, const size_t id) {
+    tuner.AddParameter(id, "WGS1", {32, 64, 128, 256, 512, 1024});
+    tuner.AddParameter(id, "WGS2", {32, 64, 128, 256, 512, 1024});
+    tuner.AddParameter(id, "VW", {1});
+  }
+
+  // Sets the constraints and local memory size
+  static void SetConstraints(cltune::Tuner &, const size_t) { }
+  static void SetLocalMemorySize(cltune::Tuner &, const size_t, const Arguments<T> &) { }
+
+  // Sets the base thread configuration
+  static std::vector<size_t> GlobalSize(const Arguments<T> &) { return {2}; }
+  static std::vector<size_t> GlobalSizeRef(const Arguments<T> &) { return {2*64*64}; }
+  static std::vector<size_t> LocalSize() { return {1}; }
+  static std::vector<size_t> LocalSizeRef() { return {64}; }
+
+  // Transforms the thread configuration based on the parameters
+  using TransformVector = std::vector<std::vector<std::string>>;
+  static TransformVector MulLocal() { return {{"WGS1"}}; }
+  static TransformVector DivLocal() { return {}; }
+  static TransformVector MulGlobal() { return {{"WGS1"},{"WGS2"}}; }
+  static TransformVector DivGlobal() { return {}; }
+
+  // Sets the kernel's arguments
+  static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args,
+                           std::vector<T> &x_vec, std::vector<T> &y_vec,
+                           std::vector<T> &, std::vector<T> &, std::vector<T> &,
+                           std::vector<T> &temp) {
+    tuner.AddArgumentScalar(static_cast<int>(args.n));
+    tuner.AddArgumentInput(x_vec);
+    tuner.AddArgumentScalar(0);
+    tuner.AddArgumentScalar(1);
+    tuner.AddArgumentInput(y_vec);
+    tuner.AddArgumentScalar(0);
+    tuner.AddArgumentScalar(1);
+    tuner.AddArgumentInput(temp); // No output checking for the result - size varies
+    tuner.AddArgumentScalar(static_cast<int>(false));
+  }
+
+  // Describes how to compute the performance metrics
+  static size_t GetMetric(const Arguments<T> &args) {
+    return (2*args.n + 1) * GetBytes(args.precision);
+  }
+  static std::string PerformanceUnit() { return "GB/s"; }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// Shortcuts to the clblast namespace
+using float2 = clblast::float2;
+using double2 = clblast::double2;
+
+// Main function (not within the clblast namespace)
+int main(int argc, char *argv[]) {
+  switch(clblast::GetPrecision(argc, argv)) {
+    case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
+    case clblast::Precision::kSingle: clblast::Tuner<clblast::TuneXdot<float>, float>(argc, argv); break;
+    case clblast::Precision::kDouble: clblast::Tuner<clblast::TuneXdot<double>, double>(argc, argv); break;
+    case clblast::Precision::kComplexSingle: clblast::Tuner<clblast::TuneXdot<float2>, float2>(argc, argv); break;
+    case clblast::Precision::kComplexDouble: clblast::Tuner<clblast::TuneXdot<double2>, double2>(argc, argv); break;
+  }
+  return 0;
+}
+
+// =================================================================================================