Merge branch 'ger_routines' into development

79 files changed, 2433 insertions, 134 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 5918d3eb..e35fda7a 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -107,11 +107,12 @@ include_directories(${clblast_SOURCE_DIR}/include ${OPENCL_INCLUDE_DIRS})
 # ==================================================================================================
 
 # Sets the supported routines and the used kernels. New routines and kernels should be added here.
-set(KERNELS copy pad transpose padtranspose xaxpy xdot xgemm xgemv)
+set(KERNELS copy pad transpose padtranspose xaxpy xdot xger xgemm xgemv)
 set(SAMPLE_PROGRAMS_CPP sgemm)
 set(SAMPLE_PROGRAMS_C sgemm)
 set(LEVEL1_ROUTINES xswap xscal xcopy xaxpy xdot xdotu xdotc)
-set(LEVEL2_ROUTINES xgemv xgbmv xhemv xhbmv xhpmv xsymv xsbmv xspmv xtrmv xtbmv xtpmv)
+set(LEVEL2_ROUTINES xgemv xgbmv xhemv xhbmv xhpmv xsymv xsbmv xspmv xtrmv xtbmv xtpmv
+                    xger xgeru xgerc xher xhpr xsyr xspr)
 set(LEVEL3_ROUTINES xgemm xsymm xhemm xsyrk xherk xsyr2k xher2k xtrmm)
 set(ROUTINES ${LEVEL1_ROUTINES} ${LEVEL2_ROUTINES} ${LEVEL3_ROUTINES})
 set(PRECISIONS 32 64 3232 6464)
diff --git a/README.md b/README.md
index 2add9798..1d6e0bad 100644
--- a/README.md
+++ b/README.md
@@ -155,7 +155,7 @@ These graphs can be generated automatically on your own device. First, compile C
 Supported routines
 -------------
 
-CLBlast is in active development but already supports the majority of BLAS routines. The currently supported routines are marked with '✔' in the following tables:
+CLBlast is in active development but already supports almost all the BLAS routines. The currently supported routines are marked with '✔' in the following tables:
 
 | Level-1  | S | D | C | Z | Notes   |
 | ---------|---|---|---|---|---------|
@@ -191,15 +191,15 @@ CLBlast is in active development but already supports the majority of BLAS routi
 | xTRSV    |   |   |   |   |         |
 | xTBSV    |   |   |   |   |         |
 | xTPSV    |   |   |   |   |         |
-| xGER     |   |   | - | - |         |
-| xGERU    | - | - |   |   |         |
-| xGERC    | - | - |   |   |         |
-| xHER     | - | - |   |   |         |
-| xHPR     | - | - |   |   |         |
+| xGER     | ✔ | ✔ | - | - |         |
+| xGERU    | - | - | ✔ | ✔ |         |
+| xGERC    | - | - | ✔ | ✔ |         |
+| xHER     | - | - | ✔ | ✔ |         |
+| xHPR     | - | - | ✔ | ✔ |         |
 | xHER2    | - | - |   |   |         |
 | xHPR2    | - | - |   |   |         |
-| xSYR     |   |   | - | - |         |
-| xSPR     |   |   | - | - |         |
+| xSYR     | ✔ | ✔ | - | - |         |
+| xSPR     | ✔ | ✔ | - | - |         |
 | xSYR2    |   |   | - | - |         |
 | xSPR2    |   |   | - | - |         |
 
diff --git a/include/internal/database.h b/include/internal/database.h
index 08e449fa..ca79fdad 100644
--- a/include/internal/database.h
+++ b/include/internal/database.h
@@ -57,22 +57,20 @@ class Database {
 
   // The OpenCL device vendors
   static constexpr auto kDeviceVendorAll = "default";
-  static constexpr auto kDeviceVendorIntel = "Intel";
-  static constexpr auto kDeviceVendorAMD = "AMD";
-  static constexpr auto kDeviceVendorNVIDIA = "NVIDIA";
 
-  // Alternative names for the above vendors
+  // Alternative names for some OpenCL vendors
   const std::unordered_map<std::string,std::string> kVendorNames {
-    {"Intel(R) Corporation", kDeviceVendorIntel},
-    {"GenuineIntel", kDeviceVendorIntel},
-    {"Advanced Micro Devices, Inc.", kDeviceVendorAMD},
-    {"NVIDIA Corporation", kDeviceVendorNVIDIA},
+    {"Intel(R) Corporation", "Intel"},
+    {"GenuineIntel", "Intel"},
+    {"Advanced Micro Devices, Inc.", "AMD"},
+    {"NVIDIA Corporation", "NVIDIA"},
   };
 
   // The database consists of separate database entries, stored together in a vector
   static const DatabaseEntry XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble;
   static const DatabaseEntry XdotSingle, XdotDouble, XdotComplexSingle, XdotComplexDouble;
   static const DatabaseEntry XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble;
+  static const DatabaseEntry XgerSingle, XgerDouble, XgerComplexSingle, XgerComplexDouble;
   static const DatabaseEntry XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble;
   static const DatabaseEntry CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble;
   static const DatabaseEntry PadSingle, PadDouble, PadComplexSingle, PadComplexDouble;
diff --git a/include/internal/database/xger.h b/include/internal/database/xger.h
new file mode 100644
index 00000000..c9cfb6cd
--- /dev/null
+++ b/include/internal/database/xger.h
@@ -0,0 +1,88 @@
+
+// =================================================================================================
+// 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):
+//   Database generator <database.py>
+//
+// This file populates the database with best-found tuning parameters for the 'Xger' kernels.
+//
+// =================================================================================================
+
+namespace clblast {
+// =================================================================================================
+
+const Database::DatabaseEntry Database::XgerSingle = {
+  "Xger", Precision::kSingle, {
+    { // Intel CPUs
+      kDeviceTypeCPU, "Intel", {
+        { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz",        { {"WGS1",128}, {"WGS2",2}, {"WPT",4} } },
+        { "default",                                         { {"WGS1",128}, {"WGS2",2}, {"WPT",4} } },
+      }
+    },
+    { // Default
+      kDeviceTypeAll, "default", {
+        { "default",                                         { {"WGS1",128}, {"WGS2",2}, {"WPT",4} } },
+      }
+    },
+  }
+};
+
+// =================================================================================================
+
+const Database::DatabaseEntry Database::XgerComplexSingle = {
+  "Xger", Precision::kComplexSingle, {
+    { // Intel CPUs
+      kDeviceTypeCPU, "Intel", {
+        { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz",        { {"WGS1",512}, {"WGS2",8}, {"WPT",2} } },
+        { "default",                                         { {"WGS1",512}, {"WGS2",8}, {"WPT",2} } },
+      }
+    },
+    { // Default
+      kDeviceTypeAll, "default", {
+        { "default",                                         { {"WGS1",512}, {"WGS2",8}, {"WPT",2} } },
+      }
+    },
+  }
+};
+
+// =================================================================================================
+
+const Database::DatabaseEntry Database::XgerDouble = {
+  "Xger", Precision::kDouble, {
+    { // Intel CPUs
+      kDeviceTypeCPU, "Intel", {
+        { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz",        { {"WGS1",512}, {"WGS2",16}, {"WPT",1} } },
+        { "default",                                         { {"WGS1",512}, {"WGS2",16}, {"WPT",1} } },
+      }
+    },
+    { // Default
+      kDeviceTypeAll, "default", {
+        { "default",                                         { {"WGS1",512}, {"WGS2",16}, {"WPT",1} } },
+      }
+    },
+  }
+};
+
+// =================================================================================================
+
+const Database::DatabaseEntry Database::XgerComplexDouble = {
+  "Xger", Precision::kComplexDouble, {
+    { // Intel CPUs
+      kDeviceTypeCPU, "Intel", {
+        { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz",        { {"WGS1",512}, {"WGS2",1}, {"WPT",1} } },
+        { "default",                                         { {"WGS1",512}, {"WGS2",1}, {"WPT",1} } },
+      }
+    },
+    { // Default
+      kDeviceTypeAll, "default", {
+        { "default",                                         { {"WGS1",512}, {"WGS2",1}, {"WPT",1} } },
+      }
+    },
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
diff --git a/include/internal/routines/level2/xger.h b/include/internal/routines/level2/xger.h
new file mode 100644
index 00000000..45ecea10
--- /dev/null
+++ b/include/internal/routines/level2/xger.h
@@ -0,0 +1,58 @@
+
+// =================================================================================================
+// 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 Xger routine. The precision is implemented using a template argument.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_ROUTINES_XGER_H_
+#define CLBLAST_ROUTINES_XGER_H_
+
+#include "internal/routine.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class Xger: public Routine<T> {
+ public:
+
+  // Members and methods from the base class
+  using Routine<T>::db_;
+  using Routine<T>::source_string_;
+  using Routine<T>::queue_;
+  using Routine<T>::GetProgramFromCache;
+  using Routine<T>::TestVectorX;
+  using Routine<T>::TestVectorY;
+  using Routine<T>::TestMatrixA;
+  using Routine<T>::RunKernel;
+  using Routine<T>::ErrorIn;
+
+  // Constructor
+  Xger(Queue &queue, Event &event, const std::string &name = "GER");
+
+  // Templated-precision implementation of the routine
+  StatusCode DoGer(const Layout layout,
+                   const size_t m, const size_t n,
+                   const T alpha,
+                   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 Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld);
+
+ private:
+  // Static variable to get the precision
+  const static Precision precision_;
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_ROUTINES_XGER_H_
+#endif
diff --git a/include/internal/routines/level2/xgerc.h b/include/internal/routines/level2/xgerc.h
new file mode 100644
index 00000000..8e515a14
--- /dev/null
+++ b/include/internal/routines/level2/xgerc.h
@@ -0,0 +1,46 @@
+
+// =================================================================================================
+// 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 Xgerc routine. The precision is implemented using a template argument.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_ROUTINES_XGERC_H_
+#define CLBLAST_ROUTINES_XGERC_H_
+
+#include "internal/routines/level2/xger.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class Xgerc: public Xger<T> {
+ public:
+
+  // Uses the regular Xger routine
+  using Xger<T>::DoGer;
+
+  // Constructor
+  Xgerc(Queue &queue, Event &event, const std::string &name = "GERC");
+
+  // Templated-precision implementation of the routine
+  StatusCode DoGerc(const Layout layout,
+                    const size_t m, const size_t n,
+                    const T alpha,
+                    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 Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld);
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_ROUTINES_XGERC_H_
+#endif
diff --git a/include/internal/routines/level2/xgeru.h b/include/internal/routines/level2/xgeru.h
new file mode 100644
index 00000000..ec485c37
--- /dev/null
+++ b/include/internal/routines/level2/xgeru.h
@@ -0,0 +1,46 @@
+
+// =================================================================================================
+// 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 Xgeru routine. The precision is implemented using a template argument.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_ROUTINES_XGERU_H_
+#define CLBLAST_ROUTINES_XGERU_H_
+
+#include "internal/routines/level2/xger.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class Xgeru: public Xger<T> {
+ public:
+
+  // Uses the regular Xger routine
+  using Xger<T>::DoGer;
+
+  // Constructor
+  Xgeru(Queue &queue, Event &event, const std::string &name = "GERU");
+
+  // Templated-precision implementation of the routine
+  StatusCode DoGeru(const Layout layout,
+                    const size_t m, const size_t n,
+                    const T alpha,
+                    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 Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld);
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_ROUTINES_XGERU_H_
+#endif
diff --git a/include/internal/routines/level2/xher.h b/include/internal/routines/level2/xher.h
new file mode 100644
index 00000000..6322265b
--- /dev/null
+++ b/include/internal/routines/level2/xher.h
@@ -0,0 +1,61 @@
+
+// =================================================================================================
+// 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 Xher routine. The precision is implemented using a template argument.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_ROUTINES_XHER_H_
+#define CLBLAST_ROUTINES_XHER_H_
+
+#include "internal/routine.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T, typename U>
+class Xher: public Routine<T> {
+ public:
+
+  // Members and methods from the base class
+  using Routine<T>::db_;
+  using Routine<T>::source_string_;
+  using Routine<T>::queue_;
+  using Routine<T>::GetProgramFromCache;
+  using Routine<T>::TestVectorX;
+  using Routine<T>::TestMatrixA;
+  using Routine<T>::TestMatrixAP;
+  using Routine<T>::RunKernel;
+  using Routine<T>::ErrorIn;
+
+  // Constructor
+  Xher(Queue &queue, Event &event, const std::string &name = "HER");
+
+  // Translates alpha of type 'U' into type 'T'
+  T GetAlpha(const U alpha);
+
+  // Templated-precision implementation of the routine
+  StatusCode DoHer(const Layout layout, const Triangle triangle,
+                   const size_t n,
+                   const U alpha,
+                   const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                   const Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld,
+                   const bool packed = false);
+
+ private:
+  // Static variable to get the precision
+  const static Precision precision_;
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_ROUTINES_XHER_H_
+#endif
diff --git a/include/internal/routines/level2/xhpr.h b/include/internal/routines/level2/xhpr.h
new file mode 100644
index 00000000..a0c3cb92
--- /dev/null
+++ b/include/internal/routines/level2/xhpr.h
@@ -0,0 +1,45 @@
+
+// =================================================================================================
+// 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 Xhpr routine. The precision is implemented using a template argument.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_ROUTINES_XHPR_H_
+#define CLBLAST_ROUTINES_XHPR_H_
+
+#include "internal/routines/level2/xher.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T, typename U>
+class Xhpr: public Xher<T,U> {
+ public:
+
+  // Uses the regular Xher routine
+  using Xher<T,U>::DoHer;
+
+  // Constructor
+  Xhpr(Queue &queue, Event &event, const std::string &name = "HPR");
+
+  // Templated-precision implementation of the routine
+  StatusCode DoHpr(const Layout layout, const Triangle triangle,
+                   const size_t n,
+                   const U alpha,
+                   const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                   const Buffer<T> &ap_buffer, const size_t ap_offset);
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_ROUTINES_XHPR_H_
+#endif
diff --git a/include/internal/routines/level2/xspr.h b/include/internal/routines/level2/xspr.h
new file mode 100644
index 00000000..5b01d2cb
--- /dev/null
+++ b/include/internal/routines/level2/xspr.h
@@ -0,0 +1,45 @@
+
+// =================================================================================================
+// 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 Xspr routine. The precision is implemented using a template argument.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_ROUTINES_XSPR_H_
+#define CLBLAST_ROUTINES_XSPR_H_
+
+#include "internal/routines/level2/xher.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class Xspr: public Xher<T,T> {
+ public:
+
+  // Uses the regular Xher routine
+  using Xher<T,T>::DoHer;
+
+  // Constructor
+  Xspr(Queue &queue, Event &event, const std::string &name = "SPR");
+
+  // Templated-precision implementation of the routine
+  StatusCode DoSpr(const Layout layout, const Triangle triangle,
+                   const size_t n,
+                   const T alpha,
+                   const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                   const Buffer<T> &ap_buffer, const size_t ap_offset);
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_ROUTINES_XSPR_H_
+#endif
diff --git a/include/internal/routines/level2/xsyr.h b/include/internal/routines/level2/xsyr.h
new file mode 100644
index 00000000..9704a881
--- /dev/null
+++ b/include/internal/routines/level2/xsyr.h
@@ -0,0 +1,45 @@
+
+// =================================================================================================
+// 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 Xsyr routine. The precision is implemented using a template argument.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_ROUTINES_XSYR_H_
+#define CLBLAST_ROUTINES_XSYR_H_
+
+#include "internal/routines/level2/xher.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class Xsyr: public Xher<T,T> {
+ public:
+
+  // Uses the regular Xher routine
+  using Xher<T,T>::DoHer;
+
+  // Constructor
+  Xsyr(Queue &queue, Event &event, const std::string &name = "SYR");
+
+  // Templated-precision implementation of the routine
+  StatusCode DoSyr(const Layout layout, const Triangle triangle,
+                   const size_t n,
+                   const T alpha,
+                   const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                   const Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld);
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_ROUTINES_XSYR_H_
+#endif
diff --git a/include/internal/utilities.h b/include/internal/utilities.h
index ed17271f..b6307a85 100644
--- a/include/internal/utilities.h
+++ b/include/internal/utilities.h
@@ -171,7 +171,8 @@ T GetArgument(const int argc, char *argv[], std::string &help,
               const std::string &option, const T default_value);
 
 // Returns the precision only
-Precision GetPrecision(const int argc, char *argv[]);
+Precision GetPrecision(const int argc, char *argv[],
+                       const Precision default_precision = Precision::kSingle);
 
 // As in "GetArgument", but now only checks whether an argument is given or not
 bool CheckArgument(const int argc, char *argv[], std::string &help, const std::string &option);
diff --git a/scripts/database/database.py b/scripts/database/database.py
index f2d47717..6f3ce85e 100644
--- a/scripts/database/database.py
+++ b/scripts/database/database.py
@@ -281,7 +281,7 @@ if len(glob.glob(glob_json)) >= 1:
 	SaveDatabase(database, file_db)
 
 # Retrieves the best performing results
-print("## Calculting the best results per device/kernel...")
+print("## Calculating the best results per device/kernel...")
 bests = GetBestResults(database)
 
 # Determines the defaults for other vendors and per vendor
diff --git a/scripts/generator/generator.py b/scripts/generator/generator.py
index 25f02861..0f5fbfa7 100644
--- a/scripts/generator/generator.py
+++ b/scripts/generator/generator.py
@@ -78,15 +78,15 @@ routines = [
   Routine(False, "2a", "tbsv",  T,  [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], False, "Solves a banded triangular system of equations"),
   Routine(False, "2a", "tpsv",  T,  [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["ap"], ["x"], [], False, "Solves a packed triangular system of equations"),
   # Level 2: matrix update
-  Routine(False, "2b", "ger",   T,  [S,D],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 matrix update"),
-  Routine(False, "2b", "geru",  T,  [C,Z],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 complex matrix update"),
-  Routine(False, "2b", "gerc",  T,  [C,Z],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 complex conjugated matrix update"),
-  Routine(False, "2b", "her",   Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], False, "Hermitian rank-1 matrix update"),
-  Routine(False, "2b", "hpr",   Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], False, "Hermitian packed rank-1 matrix update"),
+  Routine(True,  "2b", "ger",   T,  [S,D],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 matrix update"),
+  Routine(True,  "2b", "geru",  T,  [C,Z],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 complex matrix update"),
+  Routine(True,  "2b", "gerc",  T,  [C,Z],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 complex conjugated matrix update"),
+  Routine(True,  "2b", "her",   Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], False, "Hermitian rank-1 matrix update"),
+  Routine(True,  "2b", "hpr",   Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], False, "Hermitian packed rank-1 matrix update"),
   Routine(False, "2b", "her2",  T,  [C,Z],     ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], False, "Hermitian rank-2 matrix update"),
   Routine(False, "2b", "hpr2",  T,  [C,Z],     ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], False, "Hermitian packed rank-2 matrix update"),
-  Routine(False, "2b", "syr",   T,  [S,D],     ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], False, "Symmetric rank-1 matrix update"),
-  Routine(False, "2b", "spr",   T,  [S,D],     ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], False, "Symmetric packed rank-1 matrix update"),
+  Routine(True,  "2b", "syr",   T,  [S,D],     ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], False, "Symmetric rank-1 matrix update"),
+  Routine(True,  "2b", "spr",   T,  [S,D],     ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], False, "Symmetric packed rank-1 matrix update"),
   Routine(False, "2b", "syr2",  T,  [S,D],     ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], False, "Symmetric rank-2 matrix update"),
   Routine(False, "2b", "spr2",  T,  [S,D],     ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], False, "Symmetric packed rank-2 matrix update"),
 ],
@@ -103,7 +103,17 @@ routines = [
 ]]
 
 # ==================================================================================================
+# Translates an option name to a CLBlast data-type
+def PrecisionToFullName(x):
+	return {
+		'H': "Half",
+		'S': "Single",
+		'D': "Double",
+		'C': "ComplexSingle",
+		'Z': "ComplexDouble",
+	}[x]
 
+# ==================================================================================================
 # Separators for the BLAS levels
 separators = ["""
 // =================================================================================================
@@ -237,7 +247,7 @@ files = [
   path_clblast+"/src/clblast_c.cc",
   path_clblast+"/test/wrapper_clblas.h",
 ]
-header_lines = [84, 52, 80, 24, 22]
+header_lines = [84, 59, 80, 24, 22]
 footer_lines = [6, 3, 5, 2, 6]
 
 # Checks whether the command-line arguments are valid; exists otherwise
@@ -315,16 +325,10 @@ for level in [1,2,3]:
 			body += "using double2 = clblast::double2;\n\n"
 			body += "// Main function (not within the clblast namespace)\n"
 			body += "int main(int argc, char *argv[]) {\n"
-			body += "  switch(clblast::GetPrecision(argc, argv)) {\n"
+			default = PrecisionToFullName(routine.flavours[0].name)
+			body += "  switch(clblast::GetPrecision(argc, argv, clblast::Precision::k"+default+")) {\n"
 			for precision in ["H","S","D","C","Z"]:
-				enum = {
-				    'H': "Half",
-				    'S': "Single",
-				    'D': "Double",
-				    'C': "ComplexSingle",
-				    'Z': "ComplexDouble",
-				}[precision]
-				body += "    case clblast::Precision::k"+enum+":"
+				body += "    case clblast::Precision::k"+PrecisionToFullName(precision)+":"
 				found = False
 				for flavour in routine.flavours:
 					if flavour.name == precision:
diff --git a/src/clblast.cc b/src/clblast.cc
index 77999aaf..466de83e 100644
--- a/src/clblast.cc
+++ b/src/clblast.cc
@@ -38,6 +38,13 @@
 #include "internal/routines/level2/xtrmv.h"
 #include "internal/routines/level2/xtbmv.h"
 #include "internal/routines/level2/xtpmv.h"
+#include "internal/routines/level2/xger.h"
+#include "internal/routines/level2/xgeru.h"
+#include "internal/routines/level2/xgerc.h"
+#include "internal/routines/level2/xher.h"
+#include "internal/routines/level2/xhpr.h"
+#include "internal/routines/level2/xsyr.h"
+#include "internal/routines/level2/xspr.h"
 
 // BLAS level-3 includes
 #include "internal/routines/level3/xgemm.h"
@@ -835,14 +842,24 @@ template StatusCode Tpsv<double2>(const Layout, const Triangle, const Transpose,
 
 // General rank-1 matrix update: SGER/DGER
 template <typename T>
-StatusCode Ger(const Layout,
-               const size_t, const size_t,
-               const T,
-               const cl_mem, const size_t, const size_t,
-               const cl_mem, const size_t, const size_t,
-               cl_mem, const size_t, const size_t,
-               cl_command_queue*, cl_event*) {
-  return StatusCode::kNotImplemented;
+StatusCode Ger(const Layout layout,
+               const size_t m, const size_t n,
+               const T alpha,
+               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_mem a_buffer, const size_t a_offset, const size_t a_ld,
+               cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xger<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoGer(layout,
+                       m, n,
+                       alpha,
+                       Buffer<T>(x_buffer), x_offset, x_inc,
+                       Buffer<T>(y_buffer), y_offset, y_inc,
+                       Buffer<T>(a_buffer), a_offset, a_ld);
 }
 template StatusCode Ger<float>(const Layout,
                                const size_t, const size_t,
@@ -861,14 +878,24 @@ template StatusCode Ger<double>(const Layout,
 
 // General rank-1 complex matrix update: CGERU/ZGERU
 template <typename T>
-StatusCode Geru(const Layout,
-                const size_t, const size_t,
-                const T,
-                const cl_mem, const size_t, const size_t,
-                const cl_mem, const size_t, const size_t,
-                cl_mem, const size_t, const size_t,
-                cl_command_queue*, cl_event*) {
-  return StatusCode::kNotImplemented;
+StatusCode Geru(const Layout layout,
+                const size_t m, const size_t n,
+                const T alpha,
+                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_mem a_buffer, const size_t a_offset, const size_t a_ld,
+                cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xgeru<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoGeru(layout,
+                        m, n,
+                        alpha,
+                        Buffer<T>(x_buffer), x_offset, x_inc,
+                        Buffer<T>(y_buffer), y_offset, y_inc,
+                        Buffer<T>(a_buffer), a_offset, a_ld);
 }
 template StatusCode Geru<float2>(const Layout,
                                  const size_t, const size_t,
@@ -887,14 +914,24 @@ template StatusCode Geru<double2>(const Layout,
 
 // General rank-1 complex conjugated matrix update: CGERC/ZGERC
 template <typename T>
-StatusCode Gerc(const Layout,
-                const size_t, const size_t,
-                const T,
-                const cl_mem, const size_t, const size_t,
-                const cl_mem, const size_t, const size_t,
-                cl_mem, const size_t, const size_t,
-                cl_command_queue*, cl_event*) {
-  return StatusCode::kNotImplemented;
+StatusCode Gerc(const Layout layout,
+                const size_t m, const size_t n,
+                const T alpha,
+                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_mem a_buffer, const size_t a_offset, const size_t a_ld,
+                cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xgerc<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoGerc(layout,
+                        m, n,
+                        alpha,
+                        Buffer<T>(x_buffer), x_offset, x_inc,
+                        Buffer<T>(y_buffer), y_offset, y_inc,
+                        Buffer<T>(a_buffer), a_offset, a_ld);
 }
 template StatusCode Gerc<float2>(const Layout,
                                  const size_t, const size_t,
@@ -913,13 +950,22 @@ template StatusCode Gerc<double2>(const Layout,
 
 // Hermitian rank-1 matrix update: CHER/ZHER
 template <typename T>
-StatusCode Her(const Layout, const Triangle,
-               const size_t,
-               const T,
-               const cl_mem, const size_t, const size_t,
-               cl_mem, const size_t, const size_t,
-               cl_command_queue*, cl_event*) {
-  return StatusCode::kNotImplemented;
+StatusCode Her(const Layout layout, const Triangle triangle,
+               const size_t n,
+               const T alpha,
+               const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+               cl_mem a_buffer, const size_t a_offset, const size_t a_ld,
+               cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xher<std::complex<T>,T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoHer(layout, triangle,
+                       n,
+                       alpha,
+                       Buffer<std::complex<T>>(x_buffer), x_offset, x_inc,
+                       Buffer<std::complex<T>>(a_buffer), a_offset, a_ld);
 }
 template StatusCode Her<float>(const Layout, const Triangle,
                                const size_t,
@@ -936,13 +982,22 @@ template StatusCode Her<double>(const Layout, const Triangle,
 
 // Hermitian packed rank-1 matrix update: CHPR/ZHPR
 template <typename T>
-StatusCode Hpr(const Layout, const Triangle,
-               const size_t,
-               const T,
-               const cl_mem, const size_t, const size_t,
-               cl_mem, const size_t,
-               cl_command_queue*, cl_event*) {
-  return StatusCode::kNotImplemented;
+StatusCode Hpr(const Layout layout, const Triangle triangle,
+               const size_t n,
+               const T alpha,
+               const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+               cl_mem ap_buffer, const size_t ap_offset,
+               cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xhpr<std::complex<T>,T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoHpr(layout, triangle,
+                       n,
+                       alpha,
+                       Buffer<std::complex<T>>(x_buffer), x_offset, x_inc,
+                       Buffer<std::complex<T>>(ap_buffer), ap_offset);
 }
 template StatusCode Hpr<float>(const Layout, const Triangle,
                                const size_t,
@@ -1011,13 +1066,22 @@ template StatusCode Hpr2<double2>(const Layout, const Triangle,
 
 // Symmetric rank-1 matrix update: SSYR/DSYR
 template <typename T>
-StatusCode Syr(const Layout, const Triangle,
-               const size_t,
-               const T,
-               const cl_mem, const size_t, const size_t,
-               cl_mem, const size_t, const size_t,
-               cl_command_queue*, cl_event*) {
-  return StatusCode::kNotImplemented;
+StatusCode Syr(const Layout layout, const Triangle triangle,
+               const size_t n,
+               const T alpha,
+               const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+               cl_mem a_buffer, const size_t a_offset, const size_t a_ld,
+               cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xsyr<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoSyr(layout, triangle,
+                       n,
+                       alpha,
+                       Buffer<T>(x_buffer), x_offset, x_inc,
+                       Buffer<T>(a_buffer), a_offset, a_ld);
 }
 template StatusCode Syr<float>(const Layout, const Triangle,
                                const size_t,
@@ -1034,13 +1098,22 @@ template StatusCode Syr<double>(const Layout, const Triangle,
 
 // Symmetric packed rank-1 matrix update: SSPR/DSPR
 template <typename T>
-StatusCode Spr(const Layout, const Triangle,
-               const size_t,
-               const T,
-               const cl_mem, const size_t, const size_t,
-               cl_mem, const size_t,
-               cl_command_queue*, cl_event*) {
-  return StatusCode::kNotImplemented;
+StatusCode Spr(const Layout layout, const Triangle triangle,
+               const size_t n,
+               const T alpha,
+               const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+               cl_mem ap_buffer, const size_t ap_offset,
+               cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(*event);
+  auto routine = Xspr<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoSpr(layout, triangle,
+                       n,
+                       alpha,
+                       Buffer<T>(x_buffer), x_offset, x_inc,
+                       Buffer<T>(ap_buffer), ap_offset);
 }
 template StatusCode Spr<float>(const Layout, const Triangle,
                                const size_t,
diff --git a/src/database.cc b/src/database.cc
index ba0a56d9..addd85d3 100644
--- a/src/database.cc
+++ b/src/database.cc
@@ -15,6 +15,7 @@
 #include "internal/database/xaxpy.h"
 #include "internal/database/xdot.h"
 #include "internal/database/xgemv.h"
+#include "internal/database/xger.h"
 #include "internal/database/xgemm.h"
 #include "internal/database/copy.h"
 #include "internal/database/pad.h"
@@ -31,6 +32,7 @@ const std::vector<Database::DatabaseEntry> Database::database = {
   XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble,
   XdotSingle, XdotDouble, XdotComplexSingle, XdotComplexDouble,
   XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble,
+  XgerSingle, XgerDouble, XgerComplexSingle, XgerComplexDouble,
   XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble,
   CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble,
   PadSingle, PadDouble, PadComplexSingle, PadComplexDouble,
diff --git a/src/kernels/common.opencl b/src/kernels/common.opencl
index f2a2e7a7..973c123e 100644
--- a/src/kernels/common.opencl
+++ b/src/kernels/common.opencl
@@ -147,6 +147,13 @@ R"(
   #define AXPBY(e, a, b, c, d) e = a*b + c*d
 #endif
 
+// The scalar GER function
+#if PRECISION == 3232 || PRECISION == 6464
+  #define GER(e, a, b, c, d) real ab; ab.x = MulReal(a,b); ab.y = MulImag(a,b); e.x = MulReal(ab,c) + d.x; e.y = MulImag(ab,c) + d.y
+#else
+  #define GER(e, a, b, c, d) e = a*b*c + d
+#endif
+
 // The complex conjugate operation for complex transforms
 #if PRECISION == 3232 || PRECISION == 6464
   #define COMPLEX_CONJUGATE(value) value.x = value.x; value.y = -value.y
diff --git a/src/kernels/level2/level2.opencl b/src/kernels/level2/level2.opencl
new file mode 100644
index 00000000..ad92595a
--- /dev/null
+++ b/src/kernels/level2/level2.opencl
@@ -0,0 +1,102 @@
+
+// =================================================================================================
+// 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 common functions for matrix update kernels (Xger, Xher).
+//
+// =================================================================================================
+
+// 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 8    // The local work-group size in first dimension
+#endif
+#ifndef WGS2
+  #define WGS2 8    // The local work-group size in second dimension
+#endif
+#ifndef WPT
+  #define WPT 1     // The amount of work-per-thread in both dimensions
+#endif
+
+// =================================================================================================
+
+// Returns an element from a vector
+inline real LoadVector(const int id, const int max,
+                       __global real* gm, const int offset, const int inc,
+                       const int do_conjugate) {
+  if (id < max) {
+    real result = gm[id*inc + offset];
+    if (do_conjugate) {
+      #if defined(ROUTINE_GERC)
+        COMPLEX_CONJUGATE(result);
+      #endif
+      #if defined(ROUTINE_HER) || defined(ROUTINE_HPR)
+        COMPLEX_CONJUGATE(result);
+      #endif
+    }
+    return result;
+  }
+  else {
+    real default_result;
+    SetToZero(default_result);
+    return default_result;
+  }
+}
+
+// Performs the rank-1 matrix update
+inline void MatrixUpdate(const int id1, const int id2, const int max1, const int max2,
+                         __global real* agm, const int a_offset, const int a_ld,
+                         const real alpha, const real xvalue, const real yvalue,
+                         const int is_upper) {
+
+  // Bounds of a regular matrix
+  if (id1 < max1 && id2 < max2) {
+
+    #if defined(ROUTINE_SPR) || defined(ROUTINE_HPR)
+      int a_index;
+      if (is_upper) {
+        a_index = (id1 <= id2) ? ((id2+1)*id2)/2 + id1 : ((id1+1)*id1)/2 + id2;
+      }
+      else {
+        a_index = (id1 >= id2) ? ((2*a_ld-(id2+1))*id2)/2 + id1 : ((2*a_ld-(id1+1))*id1)/2 + id2;
+      }
+      a_index += a_offset;
+    #else
+      const int a_index = id2*a_ld + id1 + a_offset;
+    #endif
+
+    // Loads the current value of the A matrix
+    const real avalue = agm[a_index];
+
+    // Computes result = alpha * x[i] * y[j] + a[i][j]
+    real result;
+    GER(result, alpha, xvalue, yvalue, avalue);
+
+    // For hermetian matrices
+    #if defined(ROUTINE_HER) || defined(ROUTINE_HPR)
+      if (id1 == id2) { result.y = ZERO; }
+    #endif
+    
+    // Stores the final result
+    agm[a_index] = result;
+  }
+}
+
+// =================================================================================================
+
+// End of the C++11 raw string literal
+)"
+
+// =================================================================================================
diff --git a/src/kernels/level2/xger.opencl b/src/kernels/level2/xger.opencl
new file mode 100644
index 00000000..d377fbb0
--- /dev/null
+++ b/src/kernels/level2/xger.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 Xger kernels for rank-1 matrix update.
+//
+// =================================================================================================
+
+// 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"(
+
+// =================================================================================================
+
+// Regular version of the rank-1 matrix update kernel (GER, GERU, GERC)
+__attribute__((reqd_work_group_size(WGS1, WGS2, 1)))
+__kernel void Xger(const int max1, const int max2, const real alpha,
+                   const __global real* restrict xgm, const int x_offset, const int x_inc,
+                   const __global real* ygm, const int y_offset, const int y_inc,
+                   __global real* restrict agm, const int a_offset, const int a_ld,
+                   const int is_rowmajor) {
+
+  // Register storage for X and Y
+  real xvalues[WPT];
+  real yvalues[WPT];
+
+  // Row-major version
+  if (is_rowmajor) {
+
+    // Loads the X-vector
+    #pragma unroll
+    for (int w=0; w<WPT; ++w) {
+      const int id2 = w*get_global_size(1) + get_global_id(1);
+      xvalues[w] = LoadVector(id2, max2, xgm, x_offset, x_inc, false);
+    }
+
+    // Loads the Y-vector
+    #pragma unroll
+    for (int w=0; w<WPT; ++w) {
+      const int id1 = w*get_global_size(0) + get_global_id(0);
+      yvalues[w] = LoadVector(id1, max1, ygm, y_offset, y_inc, true);
+    }
+
+    // Loops over the work per thread twice
+    #pragma unroll
+    for (int w1=0; w1<WPT; ++w1) {
+      #pragma unroll
+      for (int w2=0; w2<WPT; ++w2) {
+
+        // Global thread IDs
+        const int id1 = w1*get_global_size(0) + get_global_id(0);
+        const int id2 = w2*get_global_size(1) + get_global_id(1);
+
+        // Loads A, performs the operation, and stores the result into A
+        MatrixUpdate(id1, id2, max1, max2, agm, a_offset, a_ld,
+                     alpha, xvalues[w2], yvalues[w1], false);
+      }
+    }
+  }
+
+  // Col-major version
+  else {
+
+    // Loads the X-vector
+    #pragma unroll
+    for (int w=0; w<WPT; ++w) {
+      const int id1 = w*get_global_size(0) + get_global_id(0);
+      xvalues[w] = LoadVector(id1, max1, xgm, x_offset, x_inc, false);
+    }
+
+    // Loads the Y-vector
+    #pragma unroll
+    for (int w=0; w<WPT; ++w) {
+      const int id2 = w*get_global_size(1) + get_global_id(1);
+      yvalues[w] = LoadVector(id2, max2, ygm, y_offset, y_inc, true);
+    }
+
+    // Loops over the work per thread twice
+    #pragma unroll
+    for (int w1=0; w1<WPT; ++w1) {
+      #pragma unroll
+      for (int w2=0; w2<WPT; ++w2) {
+
+        // Global thread IDs
+        const int id1 = w1*get_global_size(0) + get_global_id(0);
+        const int id2 = w2*get_global_size(1) + get_global_id(1);
+
+        // Loads A, performs the operation, and stores the result into A
+        MatrixUpdate(id1, id2, max1, max2, agm, a_offset, a_ld,
+                     alpha, xvalues[w1], yvalues[w2], false);
+      }
+    }
+  }
+}
+
+// =================================================================================================
+
+// End of the C++11 raw string literal
+)"
+
+// =================================================================================================
diff --git a/src/kernels/level2/xher.opencl b/src/kernels/level2/xher.opencl
new file mode 100644
index 00000000..edb94ca8
--- /dev/null
+++ b/src/kernels/level2/xher.opencl
@@ -0,0 +1,73 @@
+
+// =================================================================================================
+// 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 Xher kernels for rank-1 matrix update.
+//
+// =================================================================================================
+
+// 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"(
+
+// =================================================================================================
+
+// Symmetric version of the rank-1 matrix update kernel (HER, HPR, SYR, SPR)
+__attribute__((reqd_work_group_size(WGS1, WGS2, 1)))
+__kernel void Xher(const int n, const real alpha,
+                   const __global real* restrict xgm, const int x_offset, const int x_inc,
+                   __global real* restrict agm, const int a_offset, const int a_ld,
+                   const int is_upper, const int is_rowmajor) {
+
+  // Register storage for X and XT
+  real xvalues[WPT];
+  real xtvalues[WPT];
+
+  // Loads the X-vector
+  #pragma unroll
+  for (int w=0; w<WPT; ++w) {
+    const int id2 = w*get_global_size(1) + get_global_id(1);
+    xvalues[w] = LoadVector(id2, n, xgm, x_offset, x_inc, !is_rowmajor);
+  }
+
+  // Loads the X-transposed-vector
+  #pragma unroll
+  for (int w=0; w<WPT; ++w) {
+    const int id1 = w*get_global_size(0) + get_global_id(0);
+    xtvalues[w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor);
+  }
+
+  // Loops over the work per thread twice
+  #pragma unroll
+  for (int w1=0; w1<WPT; ++w1) {
+    #pragma unroll
+    for (int w2=0; w2<WPT; ++w2) {
+
+      // Global thread IDs
+      const int id1 = w1*get_global_size(0) + get_global_id(0);
+      const int id2 = w2*get_global_size(1) + get_global_id(1);
+
+      // Skip these threads if they do not contain threads contributing to the matrix-triangle
+      if ((is_upper && (id1 > id2)) || (!is_upper && (id2 > id1))) {
+        // Do nothing
+      }
+
+      // Loads A, performs the operation, and stores the result into A
+      else {
+        MatrixUpdate(id1, id2, n, n, agm, a_offset, a_ld, alpha, xvalues[w2], xtvalues[w1], is_upper);
+      }
+    }
+  }
+}
+
+// =================================================================================================
+
+// End of the C++11 raw string literal
+)"
+
+// =================================================================================================
diff --git a/src/routines/level1/xdotu.cc b/src/routines/level1/xdotu.cc
index 0b1bd2a8..28d9b730 100644
--- a/src/routines/level1/xdotu.cc
+++ b/src/routines/level1/xdotu.cc
@@ -14,7 +14,6 @@
 #include "internal/routines/level1/xdotu.h"
 
 #include <string>
-#include <vector>
 
 namespace clblast {
 // =================================================================================================
diff --git a/src/routines/level2/xger.cc b/src/routines/level2/xger.cc
new file mode 100644
index 00000000..0953c8bb
--- /dev/null
+++ b/src/routines/level2/xger.cc
@@ -0,0 +1,112 @@
+
+// =================================================================================================
+// 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 Xger class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level2/xger.h"
+
+#include <string>
+#include <vector>
+
+namespace clblast {
+// =================================================================================================
+
+// Specific implementations to get the memory-type based on a template argument
+template <> const Precision Xger<float>::precision_ = Precision::kSingle;
+template <> const Precision Xger<double>::precision_ = Precision::kDouble;
+template <> const Precision Xger<float2>::precision_ = Precision::kComplexSingle;
+template <> const Precision Xger<double2>::precision_ = Precision::kComplexDouble;
+
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xger<T>::Xger(Queue &queue, Event &event, const std::string &name):
+    Routine<T>(queue, event, name, {"Xger"}, precision_) {
+  source_string_ =
+    #include "../../kernels/level2/level2.opencl"
+    #include "../../kernels/level2/xger.opencl"
+  ;
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xger<T>::DoGer(const Layout layout,
+                          const size_t m, const size_t n,
+                          const T alpha,
+                          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 Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld) {
+
+  // Makes sure all dimensions are larger than zero
+  if (m == 0 || n == 0) { return StatusCode::kInvalidDimension; }
+
+  // Computes whether or not the matrix has an alternative layout (row or column-major).
+  const auto a_is_rowmajor = (layout == Layout::kRowMajor);
+  const auto a_one = (a_is_rowmajor) ? n : m;
+  const auto a_two = (a_is_rowmajor) ? m : n;
+
+  // Tests the matrix and the vectors for validity
+  auto status = TestMatrixA(a_one, a_two, a_buffer, a_offset, a_ld, sizeof(T));
+  if (ErrorIn(status)) { return status; }
+  status = TestVectorX(m, 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; }
+
+  // Retrieves the Xgemv kernel from the compiled binary
+  try {
+    auto& program = GetProgramFromCache();
+    auto kernel = Kernel(program, "Xger");
+
+    // Sets the kernel arguments
+    kernel.SetArgument(0, static_cast<int>(a_one));
+    kernel.SetArgument(1, static_cast<int>(a_two));
+    kernel.SetArgument(2, alpha);
+    kernel.SetArgument(3, x_buffer());
+    kernel.SetArgument(4, static_cast<int>(x_offset));
+    kernel.SetArgument(5, static_cast<int>(x_inc));
+    kernel.SetArgument(6, y_buffer());
+    kernel.SetArgument(7, static_cast<int>(y_offset));
+    kernel.SetArgument(8, static_cast<int>(y_inc));
+    kernel.SetArgument(9, a_buffer());
+    kernel.SetArgument(10, static_cast<int>(a_offset));
+    kernel.SetArgument(11, static_cast<int>(a_ld));
+    kernel.SetArgument(12, static_cast<int>(a_is_rowmajor));
+
+    // Launches the kernel
+    auto a_one_ceiled = CeilDiv(Ceil(a_one, db_["WGS1"]), db_["WPT"]);
+    auto a_two_ceiled = CeilDiv(Ceil(a_two, db_["WGS2"]), db_["WPT"]);
+    auto global = std::vector<size_t>{a_one_ceiled, a_two_ceiled};
+    auto local = std::vector<size_t>{db_["WGS1"], db_["WGS2"]};
+    status = RunKernel(kernel, global, local);
+    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 Xger<float>;
+template class Xger<double>;
+template class Xger<float2>;
+template class Xger<double2>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/routines/level2/xgerc.cc b/src/routines/level2/xgerc.cc
new file mode 100644
index 00000000..09408898
--- /dev/null
+++ b/src/routines/level2/xgerc.cc
@@ -0,0 +1,53 @@
+
+// =================================================================================================
+// 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 Xgerc class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level2/xgerc.h"
+
+#include <string>
+
+namespace clblast {
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xgerc<T>::Xgerc(Queue &queue, Event &event, const std::string &name):
+    Xger<T>(queue, event, name) {
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xgerc<T>::DoGerc(const Layout layout,
+                            const size_t m, const size_t n,
+                            const T alpha,
+                            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 Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld) {
+
+  // Regular Ger operation on complex data, plus conjugation in the kernel guarded by the
+  // ROUTINE_GERC guard.
+  return DoGer(layout, m, n, alpha,
+               x_buffer, x_offset, x_inc,
+               y_buffer, y_offset, y_inc,
+               a_buffer, a_offset, a_ld);
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xgerc<float2>;
+template class Xgerc<double2>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/routines/level2/xgeru.cc b/src/routines/level2/xgeru.cc
new file mode 100644
index 00000000..36fd9d0a
--- /dev/null
+++ b/src/routines/level2/xgeru.cc
@@ -0,0 +1,52 @@
+
+// =================================================================================================
+// 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 Xgeru class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level2/xgeru.h"
+
+#include <string>
+
+namespace clblast {
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xgeru<T>::Xgeru(Queue &queue, Event &event, const std::string &name):
+    Xger<T>(queue, event, name) {
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xgeru<T>::DoGeru(const Layout layout,
+                            const size_t m, const size_t n,
+                            const T alpha,
+                            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 Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld) {
+
+  // Regular Ger operation on complex data
+  return DoGer(layout, m, n, alpha,
+               x_buffer, x_offset, x_inc,
+               y_buffer, y_offset, y_inc,
+               a_buffer, a_offset, a_ld);
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xgeru<float2>;
+template class Xgeru<double2>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/routines/level2/xher.cc b/src/routines/level2/xher.cc
new file mode 100644
index 00000000..5eca44b0
--- /dev/null
+++ b/src/routines/level2/xher.cc
@@ -0,0 +1,122 @@
+
+// =================================================================================================
+// 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 Xher class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level2/xher.h"
+
+#include <string>
+
+namespace clblast {
+// =================================================================================================
+
+// Specific implementations to get the memory-type based on a template argument
+template <> const Precision Xher<float, float>::precision_ = Precision::kSingle;
+template <> const Precision Xher<double, double>::precision_ = Precision::kDouble;
+template <> const Precision Xher<float2, float>::precision_ = Precision::kComplexSingle;
+template <> const Precision Xher<double2, double>::precision_ = Precision::kComplexDouble;
+
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T, typename U>
+Xher<T,U>::Xher(Queue &queue, Event &event, const std::string &name):
+    Routine<T>(queue, event, name, {"Xger"}, precision_) {
+  source_string_ =
+    #include "../../kernels/level2/level2.opencl"
+    #include "../../kernels/level2/xher.opencl"
+  ;
+}
+
+// =================================================================================================
+
+// Specializations to compute alpha of type 'T'
+template <> float2 Xher<float2,float>::GetAlpha(const float alpha) { return float2{alpha, 0.0f}; }
+template <> double2 Xher<double2,double>::GetAlpha(const double alpha) { return double2{alpha, 0.0}; }
+template <> float Xher<float,float>::GetAlpha(const float alpha) { return alpha; }
+template <> double Xher<double,double>::GetAlpha(const double alpha) { return alpha; }
+
+// =================================================================================================
+
+// The main routine
+template <typename T, typename U>
+StatusCode Xher<T,U>::DoHer(const Layout layout, const Triangle triangle,
+                            const size_t n,
+                            const U alpha,
+                            const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                            const Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld,
+                            const bool packed) {
+
+  // Makes sure the dimensions are larger than zero
+  if (n == 0) { return StatusCode::kInvalidDimension; }
+
+  // The data is either in the upper or lower triangle
+  const auto is_upper = ((triangle == Triangle::kUpper && layout != Layout::kRowMajor) ||
+                         (triangle == Triangle::kLower && layout == Layout::kRowMajor));
+  const auto is_rowmajor = (layout == Layout::kRowMajor);
+
+  // Creates a matching version of alpha
+  const auto matching_alpha = GetAlpha(alpha);
+
+  // Tests the matrix and the vectors for validity
+  auto status = StatusCode::kSuccess;
+  if (packed) { status = TestMatrixAP(n, a_buffer, a_offset, sizeof(T)); }
+  else { status = TestMatrixA(n, n, a_buffer, a_offset, a_ld, sizeof(T)); }
+  if (ErrorIn(status)) { return status; }
+  status = TestVectorX(n, x_buffer, x_offset, x_inc, sizeof(T));
+  if (ErrorIn(status)) { return status; }
+
+  // If alpha is zero an update is not required
+  if (alpha == U{0}) { return StatusCode::kSuccess; }
+
+  // Retrieves the Xgemv kernel from the compiled binary
+  try {
+    auto& program = GetProgramFromCache();
+    auto kernel = Kernel(program, "Xher");
+
+    // Sets the kernel arguments
+    kernel.SetArgument(0, static_cast<int>(n));
+    kernel.SetArgument(1, matching_alpha);
+    kernel.SetArgument(2, x_buffer());
+    kernel.SetArgument(3, static_cast<int>(x_offset));
+    kernel.SetArgument(4, static_cast<int>(x_inc));
+    kernel.SetArgument(5, a_buffer());
+    kernel.SetArgument(6, static_cast<int>(a_offset));
+    kernel.SetArgument(7, static_cast<int>(a_ld));
+    kernel.SetArgument(8, static_cast<int>(is_upper));
+    kernel.SetArgument(9, static_cast<int>(is_rowmajor));
+
+    // Launches the kernel
+    auto global_one = CeilDiv(Ceil(n, db_["WGS1"]), db_["WPT"]);
+    auto global_two = CeilDiv(Ceil(n, db_["WGS2"]), db_["WPT"]);
+    auto global = std::vector<size_t>{global_one, global_two};
+    auto local = std::vector<size_t>{db_["WGS1"], db_["WGS2"]};
+    status = RunKernel(kernel, global, local);
+    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 Xher<float, float>;
+template class Xher<double, double>;
+template class Xher<float2, float>;
+template class Xher<double2, double>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/routines/level2/xhpr.cc b/src/routines/level2/xhpr.cc
new file mode 100644
index 00000000..abe00669
--- /dev/null
+++ b/src/routines/level2/xhpr.cc
@@ -0,0 +1,50 @@
+
+// =================================================================================================
+// 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 Xhpr class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level2/xhpr.h"
+
+#include <string>
+
+namespace clblast {
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T, typename U>
+Xhpr<T,U>::Xhpr(Queue &queue, Event &event, const std::string &name):
+    Xher<T,U>(queue, event, name) {
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T, typename U>
+StatusCode Xhpr<T,U>::DoHpr(const Layout layout, const Triangle triangle,
+                            const size_t n,
+                            const U alpha,
+                            const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                            const Buffer<T> &ap_buffer, const size_t ap_offset) {
+
+  // 
+  return DoHer(layout, triangle, n, alpha,
+               x_buffer, x_offset, x_inc,
+               ap_buffer, ap_offset, n, true);
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xhpr<float2, float>;
+template class Xhpr<double2, double>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/routines/level2/xspr.cc b/src/routines/level2/xspr.cc
new file mode 100644
index 00000000..5159ad50
--- /dev/null
+++ b/src/routines/level2/xspr.cc
@@ -0,0 +1,50 @@
+
+// =================================================================================================
+// 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 Xspr class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level2/xspr.h"
+
+#include <string>
+
+namespace clblast {
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xspr<T>::Xspr(Queue &queue, Event &event, const std::string &name):
+    Xher<T,T>(queue, event, name) {
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xspr<T>::DoSpr(const Layout layout, const Triangle triangle,
+                          const size_t n,
+                          const T alpha,
+                          const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                          const Buffer<T> &ap_buffer, const size_t ap_offset) {
+
+  // 
+  return DoHer(layout, triangle, n, alpha,
+               x_buffer, x_offset, x_inc,
+               ap_buffer, ap_offset, n, true);
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xspr<float>;
+template class Xspr<double>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/routines/level2/xsyr.cc b/src/routines/level2/xsyr.cc
new file mode 100644
index 00000000..755fde0d
--- /dev/null
+++ b/src/routines/level2/xsyr.cc
@@ -0,0 +1,50 @@
+
+// =================================================================================================
+// 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 Xsyr class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level2/xsyr.h"
+
+#include <string>
+
+namespace clblast {
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xsyr<T>::Xsyr(Queue &queue, Event &event, const std::string &name):
+    Xher<T,T>(queue, event, name) {
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xsyr<T>::DoSyr(const Layout layout, const Triangle triangle,
+                          const size_t n,
+                          const T alpha,
+                          const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc,
+                          const Buffer<T> &a_buffer, const size_t a_offset, const size_t a_ld) {
+
+  // 
+  return DoHer(layout, triangle, n, alpha,
+               x_buffer, x_offset, x_inc,
+               a_buffer, a_offset, a_ld);
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xsyr<float>;
+template class Xsyr<double>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/src/tuning/xger.cc b/src/tuning/xger.cc
new file mode 100644
index 00000000..39efdb81
--- /dev/null
+++ b/src/tuning/xger.cc
@@ -0,0 +1,129 @@
+
+// =================================================================================================
+// 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 xger OpenCL kernels.
+//
+// =================================================================================================
+
+#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 TuneXger {
+ public:
+
+  // The representative kernel and the source code
+  static std::string KernelFamily() { return "xger"; }
+  static std::string KernelName() { return "Xger"; }
+  static std::string GetSources() {
+    return
+      #include "../src/kernels/common.opencl"
+      #include "../src/kernels/level2/level2.opencl"
+      #include "../src/kernels/level2/xger.opencl"
+    ;
+  }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() { return {kArgN, kArgM, kArgAlpha}; }
+
+  // Tests for valid arguments
+  static void TestValidArguments(const Arguments<T> &) { }
+
+  // Sets the default values for the arguments
+  static size_t DefaultM() { return 1024; }
+  static size_t DefaultN() { return 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.m; }
+  static size_t GetSizeY(const Arguments<T> &args) { return args.n; }
+  static size_t GetSizeA(const Arguments<T> &args) { return args.m * args.n; }
+  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> &) { return 1; } // N/A for this kernel
+
+  // Sets the tuning parameters and their possible values
+  static void SetParameters(cltune::Tuner &tuner, const size_t id) {
+    tuner.AddParameter(id, "WGS1", {4, 8, 16, 32, 64, 128, 256, 512});
+    tuner.AddParameter(id, "WGS2", {1, 2, 4, 8, 16, 32, 64, 128, 256});
+    tuner.AddParameter(id, "WPT", {1, 2, 4});
+  }
+
+  // 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> &args) { return {args.m, args.n}; }
+  static std::vector<size_t> GlobalSizeRef(const Arguments<T> &args) { return GlobalSize(args); }
+  static std::vector<size_t> LocalSize() { return {1, 1}; }
+  static std::vector<size_t> LocalSizeRef() { return {8, 8}; }
+
+  // Transforms the thread configuration based on the parameters
+  using TransformVector = std::vector<std::vector<std::string>>;
+  static TransformVector MulLocal() { return {{"WGS1", "WGS2"}}; }
+  static TransformVector DivLocal() { return {}; }
+  static TransformVector MulGlobal() { return {}; }
+  static TransformVector DivGlobal() { return {{"WPT", "WPT"}}; }
+
+  // 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> &a_mat, std::vector<T> &, std::vector<T> &,
+                           std::vector<T> &) {
+    tuner.AddArgumentScalar(static_cast<int>(args.m));
+    tuner.AddArgumentScalar(static_cast<int>(args.n));
+    tuner.AddArgumentScalar(args.alpha);
+    tuner.AddArgumentInput(x_vec);
+    tuner.AddArgumentScalar(0); // x_offset
+    tuner.AddArgumentScalar(1); // x_increment
+    tuner.AddArgumentInput(y_vec);
+    tuner.AddArgumentScalar(0); // y_offset
+    tuner.AddArgumentScalar(1); // y_increment
+    tuner.AddArgumentOutput(a_mat);
+    tuner.AddArgumentScalar(0); // a_offset
+    tuner.AddArgumentScalar(static_cast<int>(args.m)); // a_ld
+    tuner.AddArgumentScalar(0); // a_is_rowmajor
+  }
+
+  // Describes how to compute the performance metrics
+  static size_t GetMetric(const Arguments<T> &args) {
+    return (2*args.m*args.n + args.m + args.n) * 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::TuneXger<float>, float>(argc, argv); break;
+    case clblast::Precision::kDouble: clblast::Tuner<clblast::TuneXger<double>, double>(argc, argv); break;
+    case clblast::Precision::kComplexSingle: clblast::Tuner<clblast::TuneXger<float2>, float2>(argc, argv); break;
+    case clblast::Precision::kComplexDouble: clblast::Tuner<clblast::TuneXger<double2>, double2>(argc, argv); break;
+  }
+  return 0;
+}
+
+// =================================================================================================
diff --git a/src/utilities.cc b/src/utilities.cc
index 24efb14c..68a4f02a 100644
--- a/src/utilities.cc
+++ b/src/utilities.cc
@@ -161,9 +161,9 @@ template Precision GetArgument<Precision>(const int, char **, std::string&, cons
 // =================================================================================================
 
 // Returns only the precision argument
-Precision GetPrecision(const int argc, char *argv[]) {
+Precision GetPrecision(const int argc, char *argv[], const Precision default_precision) {
   auto dummy = std::string{};
-  return GetArgument(argc, argv, dummy, kArgPrecision, Precision::kSingle);
+  return GetArgument(argc, argv, dummy, kArgPrecision, default_precision);
 }
 
 // =================================================================================================
diff --git a/test/performance/routines/level1/xaxpy.cc b/test/performance/routines/level1/xaxpy.cc
index 7ab15f28..b423bc3a 100644
--- a/test/performance/routines/level1/xaxpy.cc
+++ b/test/performance/routines/level1/xaxpy.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXaxpy<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level1/xcopy.cc b/test/performance/routines/level1/xcopy.cc
index 6277e8fb..c04c6c1c 100644
--- a/test/performance/routines/level1/xcopy.cc
+++ b/test/performance/routines/level1/xcopy.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXcopy<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level1/xdot.cc b/test/performance/routines/level1/xdot.cc
index 5aa76762..f4616464 100644
--- a/test/performance/routines/level1/xdot.cc
+++ b/test/performance/routines/level1/xdot.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXdot<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level1/xdotc.cc b/test/performance/routines/level1/xdotc.cc
index 81511085..5f36b80e 100644
--- a/test/performance/routines/level1/xdotc.cc
+++ b/test/performance/routines/level1/xdotc.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level1/xdotu.cc b/test/performance/routines/level1/xdotu.cc
index 888eede3..f19f751b 100644
--- a/test/performance/routines/level1/xdotu.cc
+++ b/test/performance/routines/level1/xdotu.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level1/xscal.cc b/test/performance/routines/level1/xscal.cc
index be49c066..bd38f43e 100644
--- a/test/performance/routines/level1/xscal.cc
+++ b/test/performance/routines/level1/xscal.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXscal<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level1/xswap.cc b/test/performance/routines/level1/xswap.cc
index 52fdc580..112641d3 100644
--- a/test/performance/routines/level1/xswap.cc
+++ b/test/performance/routines/level1/xswap.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXswap<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xgbmv.cc b/test/performance/routines/level2/xgbmv.cc
index 629e2182..b050184d 100644
--- a/test/performance/routines/level2/xgbmv.cc
+++ b/test/performance/routines/level2/xgbmv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXgbmv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xgemv.cc b/test/performance/routines/level2/xgemv.cc
index 2a1983de..51ab9a10 100644
--- a/test/performance/routines/level2/xgemv.cc
+++ b/test/performance/routines/level2/xgemv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXgemv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xger.cc b/test/performance/routines/level2/xger.cc
index 5fb0d91d..2d956346 100644
--- a/test/performance/routines/level2/xger.cc
+++ b/test/performance/routines/level2/xger.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXger<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xgerc.cc b/test/performance/routines/level2/xgerc.cc
index fd511e42..acd0fab7 100644
--- a/test/performance/routines/level2/xgerc.cc
+++ b/test/performance/routines/level2/xgerc.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xgeru.cc b/test/performance/routines/level2/xgeru.cc
index 689ab2b1..a5973777 100644
--- a/test/performance/routines/level2/xgeru.cc
+++ b/test/performance/routines/level2/xgeru.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xhbmv.cc b/test/performance/routines/level2/xhbmv.cc
index dabe6ec8..28b71045 100644
--- a/test/performance/routines/level2/xhbmv.cc
+++ b/test/performance/routines/level2/xhbmv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xhemv.cc b/test/performance/routines/level2/xhemv.cc
index 77447d76..622854a7 100644
--- a/test/performance/routines/level2/xhemv.cc
+++ b/test/performance/routines/level2/xhemv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xher.cc b/test/performance/routines/level2/xher.cc
index 4ef87e45..613d7766 100644
--- a/test/performance/routines/level2/xher.cc
+++ b/test/performance/routines/level2/xher.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xher2.cc b/test/performance/routines/level2/xher2.cc
index 2d7e17ab..c335d3be 100644
--- a/test/performance/routines/level2/xher2.cc
+++ b/test/performance/routines/level2/xher2.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xhpmv.cc b/test/performance/routines/level2/xhpmv.cc
index b9dd3f82..1e726569 100644
--- a/test/performance/routines/level2/xhpmv.cc
+++ b/test/performance/routines/level2/xhpmv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xhpr.cc b/test/performance/routines/level2/xhpr.cc
index f596682c..000b69af 100644
--- a/test/performance/routines/level2/xhpr.cc
+++ b/test/performance/routines/level2/xhpr.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xhpr2.cc b/test/performance/routines/level2/xhpr2.cc
index 1c493226..19bafc46 100644
--- a/test/performance/routines/level2/xhpr2.cc
+++ b/test/performance/routines/level2/xhpr2.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level2/xsbmv.cc b/test/performance/routines/level2/xsbmv.cc
index febc6bfd..eabab3b7 100644
--- a/test/performance/routines/level2/xsbmv.cc
+++ b/test/performance/routines/level2/xsbmv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXsbmv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xspmv.cc b/test/performance/routines/level2/xspmv.cc
index 97c6b032..2a9ef925 100644
--- a/test/performance/routines/level2/xspmv.cc
+++ b/test/performance/routines/level2/xspmv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXspmv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xspr.cc b/test/performance/routines/level2/xspr.cc
index cc18d9b6..84331d74 100644
--- a/test/performance/routines/level2/xspr.cc
+++ b/test/performance/routines/level2/xspr.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXspr<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xspr2.cc b/test/performance/routines/level2/xspr2.cc
index 768452be..c42009a1 100644
--- a/test/performance/routines/level2/xspr2.cc
+++ b/test/performance/routines/level2/xspr2.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXspr2<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xsymv.cc b/test/performance/routines/level2/xsymv.cc
index 6748026f..3f72fe77 100644
--- a/test/performance/routines/level2/xsymv.cc
+++ b/test/performance/routines/level2/xsymv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXsymv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xsyr.cc b/test/performance/routines/level2/xsyr.cc
index 84510e5d..6b31d3a9 100644
--- a/test/performance/routines/level2/xsyr.cc
+++ b/test/performance/routines/level2/xsyr.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXsyr<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xsyr2.cc b/test/performance/routines/level2/xsyr2.cc
index b8c177d8..0ad59d2d 100644
--- a/test/performance/routines/level2/xsyr2.cc
+++ b/test/performance/routines/level2/xsyr2.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXsyr2<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xtbmv.cc b/test/performance/routines/level2/xtbmv.cc
index 1663dca0..a3297f34 100644
--- a/test/performance/routines/level2/xtbmv.cc
+++ b/test/performance/routines/level2/xtbmv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXtbmv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xtbsv.cc b/test/performance/routines/level2/xtbsv.cc
index e0cb9f2e..4dcd9a06 100644
--- a/test/performance/routines/level2/xtbsv.cc
+++ b/test/performance/routines/level2/xtbsv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXtbsv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xtpmv.cc b/test/performance/routines/level2/xtpmv.cc
index 407fdc8c..72477f2d 100644
--- a/test/performance/routines/level2/xtpmv.cc
+++ b/test/performance/routines/level2/xtpmv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXtpmv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xtpsv.cc b/test/performance/routines/level2/xtpsv.cc
index e402dc60..a3e3f7f1 100644
--- a/test/performance/routines/level2/xtpsv.cc
+++ b/test/performance/routines/level2/xtpsv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXtpsv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xtrmv.cc b/test/performance/routines/level2/xtrmv.cc
index c5563240..894a7952 100644
--- a/test/performance/routines/level2/xtrmv.cc
+++ b/test/performance/routines/level2/xtrmv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXtrmv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level2/xtrsv.cc b/test/performance/routines/level2/xtrsv.cc
index 136e2108..e8c65b0f 100644
--- a/test/performance/routines/level2/xtrsv.cc
+++ b/test/performance/routines/level2/xtrsv.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXtrsv<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level3/xgemm.cc b/test/performance/routines/level3/xgemm.cc
index 2082ceac..91897ee1 100644
--- a/test/performance/routines/level3/xgemm.cc
+++ b/test/performance/routines/level3/xgemm.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXgemm<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level3/xhemm.cc b/test/performance/routines/level3/xhemm.cc
index cc68e937..87650b9e 100644
--- a/test/performance/routines/level3/xhemm.cc
+++ b/test/performance/routines/level3/xhemm.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level3/xher2k.cc b/test/performance/routines/level3/xher2k.cc
index 70d76bed..06894816 100644
--- a/test/performance/routines/level3/xher2k.cc
+++ b/test/performance/routines/level3/xher2k.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level3/xherk.cc b/test/performance/routines/level3/xherk.cc
index b3b5dddf..d6f38fb2 100644
--- a/test/performance/routines/level3/xherk.cc
+++ b/test/performance/routines/level3/xherk.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kComplexSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kDouble: throw std::runtime_error("Unsupported precision mode");
diff --git a/test/performance/routines/level3/xsymm.cc b/test/performance/routines/level3/xsymm.cc
index f2292273..e0feadd1 100644
--- a/test/performance/routines/level3/xsymm.cc
+++ b/test/performance/routines/level3/xsymm.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXsymm<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level3/xsyr2k.cc b/test/performance/routines/level3/xsyr2k.cc
index 0c8f8f7c..4a82ddc4 100644
--- a/test/performance/routines/level3/xsyr2k.cc
+++ b/test/performance/routines/level3/xsyr2k.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXsyr2k<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level3/xsyrk.cc b/test/performance/routines/level3/xsyrk.cc
index ccd4511a..70f61322 100644
--- a/test/performance/routines/level3/xsyrk.cc
+++ b/test/performance/routines/level3/xsyrk.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXsyrk<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level3/xtrmm.cc b/test/performance/routines/level3/xtrmm.cc
index 8278d077..6f6041e4 100644
--- a/test/performance/routines/level3/xtrmm.cc
+++ b/test/performance/routines/level3/xtrmm.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXtrmm<float>, float, float>(argc, argv); break;
diff --git a/test/performance/routines/level3/xtrsm.cc b/test/performance/routines/level3/xtrsm.cc
index 45f71c5e..76ef255a 100644
--- a/test/performance/routines/level3/xtrsm.cc
+++ b/test/performance/routines/level3/xtrsm.cc
@@ -18,7 +18,7 @@ using double2 = clblast::double2;
 
 // Main function (not within the clblast namespace)
 int main(int argc, char *argv[]) {
-  switch(clblast::GetPrecision(argc, argv)) {
+  switch(clblast::GetPrecision(argc, argv, clblast::Precision::kSingle)) {
     case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
     case clblast::Precision::kSingle:
       clblast::RunClient<clblast::TestXtrsm<float>, float, float>(argc, argv); break;
diff --git a/test/routines/level2/xgemv.h b/test/routines/level2/xgemv.h
index 927bfaee..2924d498 100644
--- a/test/routines/level2/xgemv.h
+++ b/test/routines/level2/xgemv.h
@@ -35,7 +35,7 @@ class TestXgemv {
   // The list of arguments relevant for this routine
   static std::vector<std::string> GetOptions() {
     return {kArgM, kArgN,
-            kArgLayout, kArgATransp, 
+            kArgLayout, kArgATransp,
             kArgALeadDim, kArgXInc, kArgYInc,
             kArgAOffset, kArgXOffset, kArgYOffset,
             kArgAlpha, kArgBeta};
diff --git a/test/routines/level2/xger.h b/test/routines/level2/xger.h
new file mode 100644
index 00000000..98296e92
--- /dev/null
+++ b/test/routines/level2/xger.h
@@ -0,0 +1,131 @@
+
+// =================================================================================================
+// 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 a class with static methods to describe the Xger routine. Examples of
+// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These
+// static methods are used by the correctness tester and the performance tester.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_TEST_ROUTINES_XGER_H_
+#define CLBLAST_TEST_ROUTINES_XGER_H_
+
+#include <vector>
+#include <string>
+
+#include "wrapper_clblas.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class TestXger {
+ public:
+
+  // The BLAS level: 1, 2, or 3
+  static size_t BLASLevel() { return 2; }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() {
+    return {kArgM, kArgN,
+            kArgLayout,
+            kArgALeadDim, kArgXInc, kArgYInc,
+            kArgAOffset, kArgXOffset, kArgYOffset,
+            kArgAlpha};
+  }
+
+  // Describes how to obtain the sizes of the buffers
+  static size_t GetSizeX(const Arguments<T> &args) {
+    return args.m * args.x_inc + args.x_offset;
+  }
+  static size_t GetSizeY(const Arguments<T> &args) {
+    return args.n * args.y_inc + args.y_offset;
+  }
+  static size_t GetSizeA(const Arguments<T> &args) {
+    auto a_rotated = (args.layout == Layout::kRowMajor);
+    auto a_two = (a_rotated) ? args.m : args.n;
+    return a_two * args.a_ld + args.a_offset;
+  }
+
+  // Describes how to set the sizes of all the buffers
+  static void SetSizes(Arguments<T> &args) {
+    args.a_size = GetSizeA(args);
+    args.x_size = GetSizeX(args);
+    args.y_size = GetSizeY(args);
+  }
+
+  // Describes what the default values of the leading dimensions of the matrices are
+  static size_t DefaultLDA(const Arguments<T> &args) { return args.n; }
+  static size_t DefaultLDB(const Arguments<T> &) { return 1; } // N/A for this routine
+  static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine
+
+  // Describes which transpose options are relevant for this routine
+  using Transposes = std::vector<Transpose>;
+  static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine
+  static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine
+
+  // Describes how to run the CLBlast routine
+  static StatusCode RunRoutine(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = Ger(args.layout,
+                      args.m, args.n, args.alpha,
+                      buffers.x_vec(), args.x_offset, args.x_inc,
+                      buffers.y_vec(), args.y_offset, args.y_inc,
+                      buffers.a_mat(), args.a_offset, args.a_ld,
+                      &queue_plain, &event);
+    clWaitForEvents(1, &event);
+    return status;
+  }
+
+  // Describes how to run the clBLAS routine (for correctness/performance comparison)
+  static StatusCode RunReference(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = clblasXger(static_cast<clblasOrder>(args.layout),
+                             args.m, args.n, args.alpha,
+                             buffers.x_vec(), args.x_offset, args.x_inc,
+                             buffers.y_vec(), args.y_offset, args.y_inc,
+                             buffers.a_mat(), args.a_offset, args.a_ld,
+                             1, &queue_plain, 0, nullptr, &event);
+    clWaitForEvents(1, &event);
+    return static_cast<StatusCode>(status);
+  }
+
+  // Describes how to download the results of the computation (more importantly: which buffer)
+  static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers<T> &buffers, Queue &queue) {
+    std::vector<T> result(args.a_size, static_cast<T>(0));
+    buffers.a_mat.Read(queue, args.a_size, result);
+    return result;
+  }
+
+  // Describes how to compute the indices of the result buffer
+  static size_t ResultID1(const Arguments<T> &args) { return args.m; }
+  static size_t ResultID2(const Arguments<T> &args) { return args.n; }
+  static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t id2) {
+    return (args.layout == Layout::kRowMajor) ?
+           id1*args.a_ld + id2 + args.a_offset:
+           id2*args.a_ld + id1 + args.a_offset;
+  }
+
+  // Describes how to compute performance metrics
+  static size_t GetFlops(const Arguments<T> &args) {
+    return 2 * args.m * args.n;
+  }
+  static size_t GetBytes(const Arguments<T> &args) {
+    return (2*args.m*args.n + args.m + args.n) * sizeof(T);
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_TEST_ROUTINES_XGER_H_
+#endif
diff --git a/test/routines/level2/xgerc.h b/test/routines/level2/xgerc.h
new file mode 100644
index 00000000..77258d92
--- /dev/null
+++ b/test/routines/level2/xgerc.h
@@ -0,0 +1,131 @@
+
+// =================================================================================================
+// 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 a class with static methods to describe the Xgerc routine. Examples of
+// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These
+// static methods are used by the correctness tester and the performance tester.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_TEST_ROUTINES_XGERC_H_
+#define CLBLAST_TEST_ROUTINES_XGERC_H_
+
+#include <vector>
+#include <string>
+
+#include "wrapper_clblas.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class TestXgerc {
+ public:
+
+  // The BLAS level: 1, 2, or 3
+  static size_t BLASLevel() { return 2; }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() {
+    return {kArgM, kArgN,
+            kArgLayout,
+            kArgALeadDim, kArgXInc, kArgYInc,
+            kArgAOffset, kArgXOffset, kArgYOffset,
+            kArgAlpha};
+  }
+
+  // Describes how to obtain the sizes of the buffers
+  static size_t GetSizeX(const Arguments<T> &args) {
+    return args.m * args.x_inc + args.x_offset;
+  }
+  static size_t GetSizeY(const Arguments<T> &args) {
+    return args.n * args.y_inc + args.y_offset;
+  }
+  static size_t GetSizeA(const Arguments<T> &args) {
+    auto a_rotated = (args.layout == Layout::kRowMajor);
+    auto a_two = (a_rotated) ? args.m : args.n;
+    return a_two * args.a_ld + args.a_offset;
+  }
+
+  // Describes how to set the sizes of all the buffers
+  static void SetSizes(Arguments<T> &args) {
+    args.a_size = GetSizeA(args);
+    args.x_size = GetSizeX(args);
+    args.y_size = GetSizeY(args);
+  }
+
+  // Describes what the default values of the leading dimensions of the matrices are
+  static size_t DefaultLDA(const Arguments<T> &args) { return args.n; }
+  static size_t DefaultLDB(const Arguments<T> &) { return 1; } // N/A for this routine
+  static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine
+
+  // Describes which transpose options are relevant for this routine
+  using Transposes = std::vector<Transpose>;
+  static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine
+  static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine
+
+  // Describes how to run the CLBlast routine
+  static StatusCode RunRoutine(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = Gerc(args.layout,
+                       args.m, args.n, args.alpha,
+                       buffers.x_vec(), args.x_offset, args.x_inc,
+                       buffers.y_vec(), args.y_offset, args.y_inc,
+                       buffers.a_mat(), args.a_offset, args.a_ld,
+                       &queue_plain, &event);
+    clWaitForEvents(1, &event);
+    return status;
+  }
+
+  // Describes how to run the clBLAS routine (for correctness/performance comparison)
+  static StatusCode RunReference(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = clblasXgerc(static_cast<clblasOrder>(args.layout),
+                              args.m, args.n, args.alpha,
+                              buffers.x_vec(), args.x_offset, args.x_inc,
+                              buffers.y_vec(), args.y_offset, args.y_inc,
+                              buffers.a_mat(), args.a_offset, args.a_ld,
+                              1, &queue_plain, 0, nullptr, &event);
+    clWaitForEvents(1, &event);
+    return static_cast<StatusCode>(status);
+  }
+
+  // Describes how to download the results of the computation (more importantly: which buffer)
+  static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers<T> &buffers, Queue &queue) {
+    std::vector<T> result(args.a_size, static_cast<T>(0));
+    buffers.a_mat.Read(queue, args.a_size, result);
+    return result;
+  }
+
+  // Describes how to compute the indices of the result buffer
+  static size_t ResultID1(const Arguments<T> &args) { return args.m; }
+  static size_t ResultID2(const Arguments<T> &args) { return args.n; }
+  static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t id2) {
+    return (args.layout == Layout::kRowMajor) ?
+           id1*args.a_ld + id2 + args.a_offset:
+           id2*args.a_ld + id1 + args.a_offset;
+  }
+
+  // Describes how to compute performance metrics
+  static size_t GetFlops(const Arguments<T> &args) {
+    return 2 * args.m * args.n;
+  }
+  static size_t GetBytes(const Arguments<T> &args) {
+    return (2*args.m*args.n + args.m + args.n) * sizeof(T);
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_TEST_ROUTINES_XGERC_H_
+#endif
diff --git a/test/routines/level2/xgeru.h b/test/routines/level2/xgeru.h
new file mode 100644
index 00000000..e5f5f235
--- /dev/null
+++ b/test/routines/level2/xgeru.h
@@ -0,0 +1,131 @@
+
+// =================================================================================================
+// 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 a class with static methods to describe the Xgeru routine. Examples of
+// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These
+// static methods are used by the correctness tester and the performance tester.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_TEST_ROUTINES_XGERU_H_
+#define CLBLAST_TEST_ROUTINES_XGERU_H_
+
+#include <vector>
+#include <string>
+
+#include "wrapper_clblas.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class TestXgeru {
+ public:
+
+  // The BLAS level: 1, 2, or 3
+  static size_t BLASLevel() { return 2; }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() {
+    return {kArgM, kArgN,
+            kArgLayout,
+            kArgALeadDim, kArgXInc, kArgYInc,
+            kArgAOffset, kArgXOffset, kArgYOffset,
+            kArgAlpha};
+  }
+
+  // Describes how to obtain the sizes of the buffers
+  static size_t GetSizeX(const Arguments<T> &args) {
+    return args.m * args.x_inc + args.x_offset;
+  }
+  static size_t GetSizeY(const Arguments<T> &args) {
+    return args.n * args.y_inc + args.y_offset;
+  }
+  static size_t GetSizeA(const Arguments<T> &args) {
+    auto a_rotated = (args.layout == Layout::kRowMajor);
+    auto a_two = (a_rotated) ? args.m : args.n;
+    return a_two * args.a_ld + args.a_offset;
+  }
+
+  // Describes how to set the sizes of all the buffers
+  static void SetSizes(Arguments<T> &args) {
+    args.a_size = GetSizeA(args);
+    args.x_size = GetSizeX(args);
+    args.y_size = GetSizeY(args);
+  }
+
+  // Describes what the default values of the leading dimensions of the matrices are
+  static size_t DefaultLDA(const Arguments<T> &args) { return args.n; }
+  static size_t DefaultLDB(const Arguments<T> &) { return 1; } // N/A for this routine
+  static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine
+
+  // Describes which transpose options are relevant for this routine
+  using Transposes = std::vector<Transpose>;
+  static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine
+  static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine
+
+  // Describes how to run the CLBlast routine
+  static StatusCode RunRoutine(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = Geru(args.layout,
+                       args.m, args.n, args.alpha,
+                       buffers.x_vec(), args.x_offset, args.x_inc,
+                       buffers.y_vec(), args.y_offset, args.y_inc,
+                       buffers.a_mat(), args.a_offset, args.a_ld,
+                       &queue_plain, &event);
+    clWaitForEvents(1, &event);
+    return status;
+  }
+
+  // Describes how to run the clBLAS routine (for correctness/performance comparison)
+  static StatusCode RunReference(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = clblasXgeru(static_cast<clblasOrder>(args.layout),
+                              args.m, args.n, args.alpha,
+                              buffers.x_vec(), args.x_offset, args.x_inc,
+                              buffers.y_vec(), args.y_offset, args.y_inc,
+                              buffers.a_mat(), args.a_offset, args.a_ld,
+                              1, &queue_plain, 0, nullptr, &event);
+    clWaitForEvents(1, &event);
+    return static_cast<StatusCode>(status);
+  }
+
+  // Describes how to download the results of the computation (more importantly: which buffer)
+  static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers<T> &buffers, Queue &queue) {
+    std::vector<T> result(args.a_size, static_cast<T>(0));
+    buffers.a_mat.Read(queue, args.a_size, result);
+    return result;
+  }
+
+  // Describes how to compute the indices of the result buffer
+  static size_t ResultID1(const Arguments<T> &args) { return args.m; }
+  static size_t ResultID2(const Arguments<T> &args) { return args.n; }
+  static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t id2) {
+    return (args.layout == Layout::kRowMajor) ?
+           id1*args.a_ld + id2 + args.a_offset:
+           id2*args.a_ld + id1 + args.a_offset;
+  }
+
+  // Describes how to compute performance metrics
+  static size_t GetFlops(const Arguments<T> &args) {
+    return 2 * args.m * args.n;
+  }
+  static size_t GetBytes(const Arguments<T> &args) {
+    return (2*args.m*args.n + args.m + args.n) * sizeof(T);
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_TEST_ROUTINES_XGERU_H_
+#endif
diff --git a/test/routines/level2/xher.h b/test/routines/level2/xher.h
new file mode 100644
index 00000000..53c4200f
--- /dev/null
+++ b/test/routines/level2/xher.h
@@ -0,0 +1,122 @@
+
+// =================================================================================================
+// 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 a class with static methods to describe the Xher routine. Examples of
+// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These
+// static methods are used by the correctness tester and the performance tester.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_TEST_ROUTINES_XHER_H_
+#define CLBLAST_TEST_ROUTINES_XHER_H_
+
+#include <vector>
+#include <string>
+
+#include "wrapper_clblas.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T, typename U>
+class TestXher {
+ public:
+
+  // The BLAS level: 1, 2, or 3
+  static size_t BLASLevel() { return 2; }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() {
+    return {kArgN,
+            kArgLayout, kArgTriangle,
+            kArgALeadDim, kArgXInc,
+            kArgAOffset, kArgXOffset,
+            kArgAlpha};
+  }
+
+  // Describes how to obtain the sizes of the buffers
+  static size_t GetSizeX(const Arguments<U> &args) {
+    return args.n * args.x_inc + args.x_offset;
+  }
+  static size_t GetSizeA(const Arguments<U> &args) {
+    return args.n * args.a_ld + args.a_offset;
+  }
+
+  // Describes how to set the sizes of all the buffers
+  static void SetSizes(Arguments<U> &args) {
+    args.a_size = GetSizeA(args);
+    args.x_size = GetSizeX(args);
+  }
+
+  // Describes what the default values of the leading dimensions of the matrices are
+  static size_t DefaultLDA(const Arguments<U> &args) { return args.n; }
+  static size_t DefaultLDB(const Arguments<U> &) { return 1; } // N/A for this routine
+  static size_t DefaultLDC(const Arguments<U> &) { return 1; } // N/A for this routine
+
+  // Describes which transpose options are relevant for this routine
+  using Transposes = std::vector<Transpose>;
+  static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine
+  static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine
+
+  // Describes how to run the CLBlast routine
+  static StatusCode RunRoutine(const Arguments<U> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = Her(args.layout, args.triangle,
+                      args.n, args.alpha,
+                      buffers.x_vec(), args.x_offset, args.x_inc,
+                      buffers.a_mat(), args.a_offset, args.a_ld,
+                      &queue_plain, &event);
+    clWaitForEvents(1, &event);
+    return status;
+  }
+
+  // Describes how to run the clBLAS routine (for correctness/performance comparison)
+  static StatusCode RunReference(const Arguments<U> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = clblasXher(static_cast<clblasOrder>(args.layout),
+                             static_cast<clblasUplo>(args.triangle),
+                             args.n, args.alpha,
+                             buffers.x_vec(), args.x_offset, args.x_inc,
+                             buffers.a_mat(), args.a_offset, args.a_ld,
+                             1, &queue_plain, 0, nullptr, &event);
+    clWaitForEvents(1, &event);
+    return static_cast<StatusCode>(status);
+  }
+
+  // Describes how to download the results of the computation (more importantly: which buffer)
+  static std::vector<T> DownloadResult(const Arguments<U> &args, Buffers<T> &buffers, Queue &queue) {
+    std::vector<T> result(args.a_size, static_cast<T>(0));
+    buffers.a_mat.Read(queue, args.a_size, result);
+    return result;
+  }
+
+  // Describes how to compute the indices of the result buffer
+  static size_t ResultID1(const Arguments<U> &args) { return args.n; }
+  static size_t ResultID2(const Arguments<U> &args) { return args.n; }
+  static size_t GetResultIndex(const Arguments<U> &args, const size_t id1, const size_t id2) {
+    return id2*args.a_ld + id1 + args.a_offset;
+  }
+
+  // Describes how to compute performance metrics
+  static size_t GetFlops(const Arguments<U> &args) {
+    return 3 * args.n * args.n;
+  }
+  static size_t GetBytes(const Arguments<U> &args) {
+    return (args.n*args.n + args.n) * sizeof(T);
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_TEST_ROUTINES_XHER_H_
+#endif
diff --git a/test/routines/level2/xhpr.h b/test/routines/level2/xhpr.h
new file mode 100644
index 00000000..03599ddc
--- /dev/null
+++ b/test/routines/level2/xhpr.h
@@ -0,0 +1,122 @@
+
+// =================================================================================================
+// 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 a class with static methods to describe the Xhpr routine. Examples of
+// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These
+// static methods are used by the correctness tester and the performance tester.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_TEST_ROUTINES_XHPR_H_
+#define CLBLAST_TEST_ROUTINES_XHPR_H_
+
+#include <vector>
+#include <string>
+
+#include "wrapper_clblas.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T, typename U>
+class TestXhpr {
+ public:
+
+  // The BLAS level: 1, 2, or 3
+  static size_t BLASLevel() { return 2; }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() {
+    return {kArgN,
+            kArgLayout, kArgTriangle,
+            kArgXInc,
+            kArgAPOffset, kArgXOffset,
+            kArgAlpha};
+  }
+
+  // Describes how to obtain the sizes of the buffers
+  static size_t GetSizeX(const Arguments<U> &args) {
+    return args.n * args.x_inc + args.x_offset;
+  }
+  static size_t GetSizeAP(const Arguments<U> &args) {
+    return ((args.n*(args.n+1)) / 2) + args.ap_offset;
+  }
+
+  // Describes how to set the sizes of all the buffers
+  static void SetSizes(Arguments<U> &args) {
+    args.ap_size = GetSizeAP(args);
+    args.x_size = GetSizeX(args);
+  }
+
+  // Describes what the default values of the leading dimensions of the matrices are
+  static size_t DefaultLDA(const Arguments<U> &args) { return args.n; }
+  static size_t DefaultLDB(const Arguments<U> &) { return 1; } // N/A for this routine
+  static size_t DefaultLDC(const Arguments<U> &) { return 1; } // N/A for this routine
+
+  // Describes which transpose options are relevant for this routine
+  using Transposes = std::vector<Transpose>;
+  static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine
+  static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine
+
+  // Describes how to run the CLBlast routine
+  static StatusCode RunRoutine(const Arguments<U> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = Hpr(args.layout, args.triangle,
+                      args.n, args.alpha,
+                      buffers.x_vec(), args.x_offset, args.x_inc,
+                      buffers.ap_mat(), args.ap_offset,
+                      &queue_plain, &event);
+    clWaitForEvents(1, &event);
+    return status;
+  }
+
+  // Describes how to run the clBLAS routine (for correctness/performance comparison)
+  static StatusCode RunReference(const Arguments<U> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = clblasXhpr(static_cast<clblasOrder>(args.layout),
+                             static_cast<clblasUplo>(args.triangle),
+                             args.n, args.alpha,
+                             buffers.x_vec(), args.x_offset, args.x_inc,
+                             buffers.ap_mat(), args.ap_offset,
+                             1, &queue_plain, 0, nullptr, &event);
+    clWaitForEvents(1, &event);
+    return static_cast<StatusCode>(status);
+  }
+
+  // Describes how to download the results of the computation (more importantly: which buffer)
+  static std::vector<T> DownloadResult(const Arguments<U> &args, Buffers<T> &buffers, Queue &queue) {
+    std::vector<T> result(args.ap_size, static_cast<T>(0));
+    buffers.ap_mat.Read(queue, args.ap_size, result);
+    return result;
+  }
+
+  // Describes how to compute the indices of the result buffer
+  static size_t ResultID1(const Arguments<U> &args) { return args.ap_size - args.ap_offset; }
+  static size_t ResultID2(const Arguments<U> &) { return 1; } // N/A for this routine
+  static size_t GetResultIndex(const Arguments<U> &args, const size_t id1, const size_t) {
+    return id1 + args.ap_offset;
+  }
+
+  // Describes how to compute performance metrics
+  static size_t GetFlops(const Arguments<U> &args) {
+    return 3 * ((args.n*(args.n+1)) / 2);
+  }
+  static size_t GetBytes(const Arguments<U> &args) {
+    return ((args.n*(args.n+1)) + args.n) * sizeof(T);
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_TEST_ROUTINES_XHPR_H_
+#endif
diff --git a/test/routines/level2/xspr.h b/test/routines/level2/xspr.h
new file mode 100644
index 00000000..819b1ca8
--- /dev/null
+++ b/test/routines/level2/xspr.h
@@ -0,0 +1,122 @@
+
+// =================================================================================================
+// 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 a class with static methods to describe the Xspr routine. Examples of
+// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These
+// static methods are used by the correctness tester and the performance tester.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_TEST_ROUTINES_XSPR_H_
+#define CLBLAST_TEST_ROUTINES_XSPR_H_
+
+#include <vector>
+#include <string>
+
+#include "wrapper_clblas.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class TestXspr {
+ public:
+
+  // The BLAS level: 1, 2, or 3
+  static size_t BLASLevel() { return 2; }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() {
+    return {kArgN,
+            kArgLayout, kArgTriangle,
+            kArgXInc,
+            kArgAPOffset, kArgXOffset,
+            kArgAlpha};
+  }
+
+  // Describes how to obtain the sizes of the buffers
+  static size_t GetSizeX(const Arguments<T> &args) {
+    return args.n * args.x_inc + args.x_offset;
+  }
+  static size_t GetSizeAP(const Arguments<T> &args) {
+    return ((args.n*(args.n+1)) / 2) + args.ap_offset;
+  }
+
+  // Describes how to set the sizes of all the buffers
+  static void SetSizes(Arguments<T> &args) {
+    args.ap_size = GetSizeAP(args);
+    args.x_size = GetSizeX(args);
+  }
+
+  // Describes what the default values of the leading dimensions of the matrices are
+  static size_t DefaultLDA(const Arguments<T> &args) { return args.n; }
+  static size_t DefaultLDB(const Arguments<T> &) { return 1; } // N/A for this routine
+  static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine
+
+  // Describes which transpose options are relevant for this routine
+  using Transposes = std::vector<Transpose>;
+  static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine
+  static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine
+
+  // Describes how to run the CLBlast routine
+  static StatusCode RunRoutine(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = Spr(args.layout, args.triangle,
+                      args.n, args.alpha,
+                      buffers.x_vec(), args.x_offset, args.x_inc,
+                      buffers.ap_mat(), args.ap_offset,
+                      &queue_plain, &event);
+    clWaitForEvents(1, &event);
+    return status;
+  }
+
+  // Describes how to run the clBLAS routine (for correctness/performance comparison)
+  static StatusCode RunReference(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = clblasXspr(static_cast<clblasOrder>(args.layout),
+                             static_cast<clblasUplo>(args.triangle),
+                             args.n, args.alpha,
+                             buffers.x_vec(), args.x_offset, args.x_inc,
+                             buffers.ap_mat(), args.ap_offset,
+                             1, &queue_plain, 0, nullptr, &event);
+    clWaitForEvents(1, &event);
+    return static_cast<StatusCode>(status);
+  }
+
+  // Describes how to download the results of the computation (more importantly: which buffer)
+  static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers<T> &buffers, Queue &queue) {
+    std::vector<T> result(args.ap_size, static_cast<T>(0));
+    buffers.ap_mat.Read(queue, args.ap_size, result);
+    return result;
+  }
+
+  // Describes how to compute the indices of the result buffer
+  static size_t ResultID1(const Arguments<T> &args) { return args.ap_size - args.ap_offset; }
+  static size_t ResultID2(const Arguments<T> &) { return 1; } // N/A for this routine
+  static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t) {
+    return id1 + args.ap_offset;
+  }
+
+  // Describes how to compute performance metrics
+  static size_t GetFlops(const Arguments<T> &args) {
+    return 3 * ((args.n*(args.n+1)) / 2);
+  }
+  static size_t GetBytes(const Arguments<T> &args) {
+    return ((args.n*(args.n+1)) + args.n) * sizeof(T);
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_TEST_ROUTINES_XSPR_H_
+#endif
diff --git a/test/routines/level2/xsyr.h b/test/routines/level2/xsyr.h
new file mode 100644
index 00000000..66b75c0c
--- /dev/null
+++ b/test/routines/level2/xsyr.h
@@ -0,0 +1,122 @@
+
+// =================================================================================================
+// 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 a class with static methods to describe the Xsyr routine. Examples of
+// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These
+// static methods are used by the correctness tester and the performance tester.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_TEST_ROUTINES_XSYR_H_
+#define CLBLAST_TEST_ROUTINES_XSYR_H_
+
+#include <vector>
+#include <string>
+
+#include "wrapper_clblas.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class TestXsyr {
+ public:
+
+  // The BLAS level: 1, 2, or 3
+  static size_t BLASLevel() { return 2; }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() {
+    return {kArgN,
+            kArgLayout, kArgTriangle,
+            kArgALeadDim, kArgXInc,
+            kArgAOffset, kArgXOffset,
+            kArgAlpha};
+  }
+
+  // Describes how to obtain the sizes of the buffers
+  static size_t GetSizeX(const Arguments<T> &args) {
+    return args.n * args.x_inc + args.x_offset;
+  }
+  static size_t GetSizeA(const Arguments<T> &args) {
+    return args.n * args.a_ld + args.a_offset;
+  }
+
+  // Describes how to set the sizes of all the buffers
+  static void SetSizes(Arguments<T> &args) {
+    args.a_size = GetSizeA(args);
+    args.x_size = GetSizeX(args);
+  }
+
+  // Describes what the default values of the leading dimensions of the matrices are
+  static size_t DefaultLDA(const Arguments<T> &args) { return args.n; }
+  static size_t DefaultLDB(const Arguments<T> &) { return 1; } // N/A for this routine
+  static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine
+
+  // Describes which transpose options are relevant for this routine
+  using Transposes = std::vector<Transpose>;
+  static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine
+  static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine
+
+  // Describes how to run the CLBlast routine
+  static StatusCode RunRoutine(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = Syr(args.layout, args.triangle,
+                      args.n, args.alpha,
+                      buffers.x_vec(), args.x_offset, args.x_inc,
+                      buffers.a_mat(), args.a_offset, args.a_ld,
+                      &queue_plain, &event);
+    clWaitForEvents(1, &event);
+    return status;
+  }
+
+  // Describes how to run the clBLAS routine (for correctness/performance comparison)
+  static StatusCode RunReference(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) {
+    auto queue_plain = queue();
+    auto event = cl_event{};
+    auto status = clblasXsyr(static_cast<clblasOrder>(args.layout),
+                             static_cast<clblasUplo>(args.triangle),
+                             args.n, args.alpha,
+                             buffers.x_vec(), args.x_offset, args.x_inc,
+                             buffers.a_mat(), args.a_offset, args.a_ld,
+                             1, &queue_plain, 0, nullptr, &event);
+    clWaitForEvents(1, &event);
+    return static_cast<StatusCode>(status);
+  }
+
+  // Describes how to download the results of the computation (more importantly: which buffer)
+  static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers<T> &buffers, Queue &queue) {
+    std::vector<T> result(args.a_size, static_cast<T>(0));
+    buffers.a_mat.Read(queue, args.a_size, result);
+    return result;
+  }
+
+  // Describes how to compute the indices of the result buffer
+  static size_t ResultID1(const Arguments<T> &args) { return args.n; }
+  static size_t ResultID2(const Arguments<T> &args) { return args.n; }
+  static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t id2) {
+    return id2*args.a_ld + id1 + args.a_offset;
+  }
+
+  // Describes how to compute performance metrics
+  static size_t GetFlops(const Arguments<T> &args) {
+    return 3 * args.n * args.n;
+  }
+  static size_t GetBytes(const Arguments<T> &args) {
+    return (args.n*args.n + args.n) * sizeof(T);
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_TEST_ROUTINES_XSYR_H_
+#endif