Merge branch 'level1_routines' into development

22 files changed, 751 insertions, 122 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index d3ad6889..8316a49a 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -118,7 +118,7 @@ include_directories(${clblast_SOURCE_DIR}/include ${OPENCL_INCLUDE_DIRS})
 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(LEVEL1_ROUTINES xswap xscal xcopy xaxpy xdot xdotu xdotc xnrm2)
 set(LEVEL2_ROUTINES xgemv xgbmv xhemv xhbmv xhpmv xsymv xsbmv xspmv xtrmv xtbmv xtpmv
                     xger xgeru xgerc xher xhpr xher2 xhpr2 xsyr xspr xsyr2 xspr2)
 set(LEVEL3_ROUTINES xgemm xsymm xhemm xsyrk xherk xsyr2k xher2k xtrmm)
diff --git a/include/clblast.h b/include/clblast.h
index ac16188f..5e5c5a46 100644
--- a/include/clblast.h
+++ b/include/clblast.h
@@ -174,6 +174,13 @@ StatusCode Dotc(const size_t n,
                 const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
                 cl_command_queue* queue, cl_event* event = nullptr);
 
+// Euclidian norm of a vector: SNRM2/DNRM2/ScNRM2/DzNRM2
+template <typename T>
+StatusCode Nrm2(const size_t n,
+                cl_mem nrm2_buffer, const size_t nrm2_offset,
+                const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                cl_command_queue* queue, cl_event* event = nullptr);
+
 // =================================================================================================
 // BLAS level-2 (matrix-vector) routines
 // =================================================================================================
diff --git a/include/clblast_c.h b/include/clblast_c.h
index a5563951..dcb3ae3a 100644
--- a/include/clblast_c.h
+++ b/include/clblast_c.h
@@ -260,6 +260,24 @@ StatusCode PUBLIC_API CLBlastZdotc(const size_t n,
                                    const cl_mem y_buffer, const size_t y_offset, const size_t y_inc,
                                    cl_command_queue* queue, cl_event* event);
 
+// Euclidian norm of a vector: SNRM2/DNRM2/ScNRM2/DzNRM2
+StatusCode PUBLIC_API CLBlastSnrm2(const size_t n,
+                                   cl_mem nrm2_buffer, const size_t nrm2_offset,
+                                   const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                                   cl_command_queue* queue, cl_event* event);
+StatusCode PUBLIC_API CLBlastDnrm2(const size_t n,
+                                   cl_mem nrm2_buffer, const size_t nrm2_offset,
+                                   const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                                   cl_command_queue* queue, cl_event* event);
+StatusCode PUBLIC_API CLBlastScnrm2(const size_t n,
+                                   cl_mem nrm2_buffer, const size_t nrm2_offset,
+                                   const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                                   cl_command_queue* queue, cl_event* event);
+StatusCode PUBLIC_API CLBlastDznrm2(const size_t n,
+                                   cl_mem nrm2_buffer, const size_t nrm2_offset,
+                                   const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                                   cl_command_queue* queue, cl_event* event);
+
 // =================================================================================================
 // BLAS level-2 (matrix-vector) routines
 // =================================================================================================
diff --git a/include/internal/routines/level1/xnrm2.h b/include/internal/routines/level1/xnrm2.h
new file mode 100644
index 00000000..b3fffef6
--- /dev/null
+++ b/include/internal/routines/level1/xnrm2.h
@@ -0,0 +1,55 @@
+
+// =================================================================================================
+// 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 Xnrm2 routine. The precision is implemented using a template argument.
+//
+// =================================================================================================
+
+#ifndef CLBLAST_ROUTINES_XNRM2_H_
+#define CLBLAST_ROUTINES_XNRM2_H_
+
+#include "internal/routine.h"
+
+namespace clblast {
+// =================================================================================================
+
+// See comment at top of file for a description of the class
+template <typename T>
+class Xnrm2: 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>::context_;
+  using Routine<T>::GetProgramFromCache;
+  using Routine<T>::TestVectorX;
+  using Routine<T>::TestVectorDot;
+  using Routine<T>::RunKernel;
+  using Routine<T>::ErrorIn;
+
+  // Constructor
+  Xnrm2(Queue &queue, Event &event, const std::string &name = "NRM2");
+
+  // Templated-precision implementation of the routine
+  StatusCode DoNrm2(const size_t n,
+                    const Buffer<T> &nrm2_buffer, const size_t nrm2_offset,
+                    const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc);
+
+ private:
+  // Static variable to get the precision
+  const static Precision precision_;
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_ROUTINES_XNRM2_H_
+#endif
diff --git a/include/internal/utilities.h b/include/internal/utilities.h
index b6307a85..35f76722 100644
--- a/include/internal/utilities.h
+++ b/include/internal/utilities.h
@@ -61,6 +61,7 @@ constexpr auto kArgBOffset = "offb";
 constexpr auto kArgCOffset = "offc";
 constexpr auto kArgAPOffset = "offap";
 constexpr auto kArgDotOffset = "offdot";
+constexpr auto kArgNrm2Offset = "offnrm2";
 constexpr auto kArgAlpha = "alpha";
 constexpr auto kArgBeta = "beta";
 
@@ -113,6 +114,7 @@ struct Arguments {
   size_t c_offset = 0;
   size_t ap_offset = 0;
   size_t dot_offset = 0;
+  size_t nrm2_offset = 0;
   T alpha = T{1.0};
   T beta = T{1.0};
   size_t x_size = 1;
@@ -121,7 +123,7 @@ struct Arguments {
   size_t b_size = 1;
   size_t c_size = 1;
   size_t ap_size = 1;
-  size_t dot_size = 1;
+  size_t scalar_size = 1;
   // Tuner-specific arguments
   double fraction = 1.0;
   // Client-specific arguments
@@ -149,7 +151,7 @@ struct Buffers {
   Buffer<T> b_mat;
   Buffer<T> c_mat;
   Buffer<T> ap_mat;
-  Buffer<T> dot;
+  Buffer<T> scalar;
 };
 
 // =================================================================================================
diff --git a/scripts/generator/datatype.py b/scripts/generator/datatype.py
index 0aa27197..9323bc4d 100644
--- a/scripts/generator/datatype.py
+++ b/scripts/generator/datatype.py
@@ -22,7 +22,8 @@ D2CL = "cl_double2"
 
 # Structure holding data-type and precision information
 class DataType():
-	def __init__(self, name, template, scalars, buffertype):
+	def __init__(self, precision_name, name, template, scalars, buffertype):
+		self.precision_name = precision_name
 		self.name = name
 		self.template = template
 		self.alpha_cpp = scalars[0]
diff --git a/scripts/generator/generator.py b/scripts/generator/generator.py
index 1eada753..6e2b2ed2 100644
--- a/scripts/generator/generator.py
+++ b/scripts/generator/generator.py
@@ -31,77 +31,80 @@ from datatype import DataType, FLT, DBL, FLT2, DBL2, F2CL, D2CL
 # ==================================================================================================
 
 # Regular data-types
-S = DataType("S", FLT,  [FLT,  FLT,  FLT,  FLT],  FLT ) # single (32)
-D = DataType("D", DBL,  [DBL,  DBL,  DBL,  DBL],  DBL ) # double (64)
-C = DataType("C", FLT2, [FLT2, FLT2, F2CL, F2CL], FLT2) # single-complex (3232)
-Z = DataType("Z", DBL2, [DBL2, DBL2, D2CL, D2CL], DBL2) # double-complex (6464)
+S = DataType("S", "S", FLT,  [FLT,  FLT,  FLT,  FLT],  FLT ) # single (32)
+D = DataType("D", "D", DBL,  [DBL,  DBL,  DBL,  DBL],  DBL ) # double (64)
+C = DataType("C", "C", FLT2, [FLT2, FLT2, F2CL, F2CL], FLT2) # single-complex (3232)
+Z = DataType("Z", "Z", DBL2, [DBL2, DBL2, D2CL, D2CL], DBL2) # double-complex (6464)
 
 # Special cases
-Css = DataType("C", FLT,          [FLT,  FLT,  FLT,  FLT], FLT2) # As C, but with constants from S
-Zdd = DataType("Z", DBL,          [DBL,  DBL,  DBL,  DBL], DBL2) # As Z, but with constants from D
-Ccs = DataType("C", FLT2+","+FLT, [FLT2, FLT,  F2CL, FLT], FLT2) # As C, but with one constant from S
-Zzd = DataType("Z", DBL2+","+DBL, [DBL2, DBL,  D2CL, DBL], DBL2) # As Z, but with one constant from D
+Sc = DataType("C", "Sc", FLT2,         [FLT2, FLT2, FLT2, FLT2], FLT2) # As C, but with real output
+Dz = DataType("Z", "Dz", DBL2,         [DBL2, DBL2, DBL2, DBL2], DBL2) # As Z, but with real output
+Css = DataType("C", "C", FLT,          [FLT,  FLT,  FLT,  FLT], FLT2) # As C, but with constants from S
+Zdd = DataType("Z", "Z", DBL,          [DBL,  DBL,  DBL,  DBL], DBL2) # As Z, but with constants from D
+Ccs = DataType("C", "C", FLT2+","+FLT, [FLT2, FLT,  F2CL, FLT], FLT2) # As C, but with one constant from S
+Zzd = DataType("Z", "Z", DBL2+","+DBL, [DBL2, DBL,  D2CL, DBL], DBL2) # As Z, but with one constant from D
 
 # C++ template data-types
-T = DataType("typename T", "T", ["T", "T", "T", "T"], "T") # regular routine
-Tc = DataType("typename T", "std::complex<T>,T", ["T", "T", "T", "T"], "std::complex<T>") # for herk
-TU = DataType("typename T, typename U", "T,U", ["T", "U", "T", "U"], "T") # for her2k
+T = DataType("T", "typename T", "T", ["T", "T", "T", "T"], "T") # regular routine
+Tc = DataType("Tc", "typename T", "std::complex<T>,T", ["T", "T", "T", "T"], "std::complex<T>") # for herk
+TU = DataType("TU", "typename T, typename U", "T,U", ["T", "U", "T", "U"], "T") # for her2k
 
 # ==================================================================================================
 
 # Populates a list of routines
 routines = [
 [ # Level 1: vector-vector
-  Routine(False, "1", "rotg",  T,  [S,D],     [], [], [], ["sa","sb","sc","ss"], [], False, "Generate givens plane rotation"),
-  Routine(False, "1", "rotmg", T,  [S,D],     [], [], [], ["sd1","sd2","sx1","sy1","sparam"], [], False, "Generate modified givens plane rotation"),
-  Routine(False, "1", "rot",   T,  [S,D],     ["n"], [], [], ["x","y"], ["cos","sin"], False, "Apply givens plane rotation"),
-  Routine(False, "1", "rotm",  T,  [S,D],     ["n"], [], [], ["x","y","sparam"], [], False, "Apply modified givens plane rotation"),
-  Routine(True,  "1", "swap",  T,  [S,D,C,Z], ["n"], [], [], ["x","y"], [], False, "Swap two vectors"),
-  Routine(True,  "1", "scal",  T,  [S,D,C,Z], ["n"], [], [], ["x"], ["alpha"], False, "Vector scaling"),
-  Routine(True,  "1", "copy",  T,  [S,D,C,Z], ["n"], [], ["x"], ["y"], [], False, "Vector copy"),
-  Routine(True,  "1", "axpy",  T,  [S,D,C,Z], ["n"], [], ["x"], ["y"], ["alpha"], False, "Vector-times-constant plus vector"),
-  Routine(True,  "1", "dot",   T,  [S,D],     ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two vectors"),
-  Routine(True,  "1", "dotu",  T,  [C,Z],     ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two complex vectors"),
-  Routine(True,  "1", "dotc",  T,  [C,Z],     ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two complex vectors, one conjugated"),
+  Routine(False, "1", "rotg",  T,  [S,D],     [], [], [], ["sa","sb","sc","ss"], [], "", "Generate givens plane rotation"),
+  Routine(False, "1", "rotmg", T,  [S,D],     [], [], [], ["sd1","sd2","sx1","sy1","sparam"], [], "", "Generate modified givens plane rotation"),
+  Routine(False, "1", "rot",   T,  [S,D],     ["n"], [], [], ["x","y"], ["cos","sin"], "", "Apply givens plane rotation"),
+  Routine(False, "1", "rotm",  T,  [S,D],     ["n"], [], [], ["x","y","sparam"], [], "", "Apply modified givens plane rotation"),
+  Routine(True,  "1", "swap",  T,  [S,D,C,Z], ["n"], [], [], ["x","y"], [], "", "Swap two vectors"),
+  Routine(True,  "1", "scal",  T,  [S,D,C,Z], ["n"], [], [], ["x"], ["alpha"], "", "Vector scaling"),
+  Routine(True,  "1", "copy",  T,  [S,D,C,Z], ["n"], [], ["x"], ["y"], [], "", "Vector copy"),
+  Routine(True,  "1", "axpy",  T,  [S,D,C,Z], ["n"], [], ["x"], ["y"], ["alpha"], "", "Vector-times-constant plus vector"),
+  Routine(True,  "1", "dot",   T,  [S,D],     ["n"], [], ["x","y"], ["dot"], [], "n", "Dot product of two vectors"),
+  Routine(True,  "1", "dotu",  T,  [C,Z],     ["n"], [], ["x","y"], ["dot"], [], "n", "Dot product of two complex vectors"),
+  Routine(True,  "1", "dotc",  T,  [C,Z],     ["n"], [], ["x","y"], ["dot"], [], "n", "Dot product of two complex vectors, one conjugated"),
+  Routine(True,  "1", "nrm2",  T, [S,D,Sc,Dz],["n"], [], ["x"], ["nrm2"], [], "2*n", "Euclidian norm of a vector"),
 ],
 [ # Level 2: matrix-vector
-  Routine(True,  "2a", "gemv",  T,  [S,D,C,Z], ["m","n"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], False, "General matrix-vector multiplication"),
-  Routine(True,  "2a", "gbmv",  T,  [S,D,C,Z], ["m","n","kl","ku"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], False, "General banded matrix-vector multiplication"),
-  Routine(True,  "2a", "hemv",  T,  [C,Z],     ["n"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], False, "Hermitian matrix-vector multiplication"),
-  Routine(True,  "2a", "hbmv",  T,  [C,Z],     ["n","k"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], False, "Hermitian banded matrix-vector multiplication"),
-  Routine(True,  "2a", "hpmv",  T,  [C,Z],     ["n"], ["layout","triangle"], ["ap","x"], ["y"], ["alpha","beta"], False, "Hermitian packed matrix-vector multiplication"),
-  Routine(True,  "2a", "symv",  T,  [S,D],     ["n"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], False, "Symmetric matrix-vector multiplication"),
-  Routine(True,  "2a", "sbmv",  T,  [S,D],     ["n","k"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], False, "Symmetric banded matrix-vector multiplication"),
-  Routine(True,  "2a", "spmv",  T,  [S,D],     ["n"], ["layout","triangle"], ["ap","x"], ["y"], ["alpha","beta"], False, "Symmetric packed matrix-vector multiplication"),
-  Routine(True,  "2a", "trmv",  T,  [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], True, "Triangular matrix-vector multiplication"),
-  Routine(True,  "2a", "tbmv",  T,  [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], True, "Triangular banded matrix-vector multiplication"),
-  Routine(True,  "2a", "tpmv",  T,  [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["ap"], ["x"], [], True, "Triangular packed matrix-vector multiplication"),
-  Routine(False, "2a", "trsv",  T,  [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], False, "Solves a triangular system of equations"),
-  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"),
+  Routine(True,  "2a", "gemv",  T,  [S,D,C,Z], ["m","n"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], "", "General matrix-vector multiplication"),
+  Routine(True,  "2a", "gbmv",  T,  [S,D,C,Z], ["m","n","kl","ku"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], "", "General banded matrix-vector multiplication"),
+  Routine(True,  "2a", "hemv",  T,  [C,Z],     ["n"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], "", "Hermitian matrix-vector multiplication"),
+  Routine(True,  "2a", "hbmv",  T,  [C,Z],     ["n","k"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], "", "Hermitian banded matrix-vector multiplication"),
+  Routine(True,  "2a", "hpmv",  T,  [C,Z],     ["n"], ["layout","triangle"], ["ap","x"], ["y"], ["alpha","beta"], "", "Hermitian packed matrix-vector multiplication"),
+  Routine(True,  "2a", "symv",  T,  [S,D],     ["n"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], "", "Symmetric matrix-vector multiplication"),
+  Routine(True,  "2a", "sbmv",  T,  [S,D],     ["n","k"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], "", "Symmetric banded matrix-vector multiplication"),
+  Routine(True,  "2a", "spmv",  T,  [S,D],     ["n"], ["layout","triangle"], ["ap","x"], ["y"], ["alpha","beta"], "", "Symmetric packed matrix-vector multiplication"),
+  Routine(True,  "2a", "trmv",  T,  [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], "n", "Triangular matrix-vector multiplication"),
+  Routine(True,  "2a", "tbmv",  T,  [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], "n", "Triangular banded matrix-vector multiplication"),
+  Routine(True,  "2a", "tpmv",  T,  [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["ap"], ["x"], [], "n", "Triangular packed matrix-vector multiplication"),
+  Routine(False, "2a", "trsv",  T,  [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], "", "Solves a triangular system of equations"),
+  Routine(False, "2a", "tbsv",  T,  [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], "", "Solves a banded triangular system of equations"),
+  Routine(False, "2a", "tpsv",  T,  [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["ap"], ["x"], [], "", "Solves a packed triangular system of equations"),
   # Level 2: 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(True,  "2b", "her2",  T,  [C,Z],     ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], False, "Hermitian rank-2 matrix update"),
-  Routine(True,  "2b", "hpr2",  T,  [C,Z],     ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], False, "Hermitian packed rank-2 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(True,  "2b", "syr2",  T,  [S,D],     ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], False, "Symmetric rank-2 matrix update"),
-  Routine(True,  "2b", "spr2",  T,  [S,D],     ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], False, "Symmetric packed rank-2 matrix update"),
+  Routine(True,  "2b", "ger",   T,  [S,D],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], "", "General rank-1 matrix update"),
+  Routine(True,  "2b", "geru",  T,  [C,Z],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], "", "General rank-1 complex matrix update"),
+  Routine(True,  "2b", "gerc",  T,  [C,Z],     ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], "", "General rank-1 complex conjugated matrix update"),
+  Routine(True,  "2b", "her",   Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], "", "Hermitian rank-1 matrix update"),
+  Routine(True,  "2b", "hpr",   Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], "", "Hermitian packed rank-1 matrix update"),
+  Routine(True,  "2b", "her2",  T,  [C,Z],     ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], "", "Hermitian rank-2 matrix update"),
+  Routine(True,  "2b", "hpr2",  T,  [C,Z],     ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], "", "Hermitian packed rank-2 matrix update"),
+  Routine(True,  "2b", "syr",   T,  [S,D],     ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], "", "Symmetric rank-1 matrix update"),
+  Routine(True,  "2b", "spr",   T,  [S,D],     ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], "", "Symmetric packed rank-1 matrix update"),
+  Routine(True,  "2b", "syr2",  T,  [S,D],     ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], "", "Symmetric rank-2 matrix update"),
+  Routine(True,  "2b", "spr2",  T,  [S,D],     ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], "", "Symmetric packed rank-2 matrix update"),
 ],
 [ # Level 3: matrix-matrix
-  Routine(True,  "3", "gemm",  T,  [S,D,C,Z], ["m","n","k"], ["layout","a_transpose","b_transpose"], ["a","b"], ["c"], ["alpha","beta"], False, "General matrix-matrix multiplication"),
-  Routine(True,  "3", "symm",  T,  [S,D,C,Z], ["m","n"], ["layout","side","triangle"], ["a","b"], ["c"], ["alpha","beta"], False, "Symmetric matrix-matrix multiplication"),
-  Routine(True,  "3", "hemm",  T,  [C,Z],     ["m","n"], ["layout","side","triangle"], ["a","b"], ["c"], ["alpha","beta"], False, "Hermitian matrix-matrix multiplication"),
-  Routine(True,  "3", "syrk",  T,  [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose"], ["a"], ["c"], ["alpha","beta"], False, "Rank-K update of a symmetric matrix"),
-  Routine(True,  "3", "herk",  Tc, [Css,Zdd], ["n","k"], ["layout","triangle","a_transpose"], ["a"], ["c"], ["alpha","beta"], False, "Rank-K update of a hermitian matrix"),
-  Routine(True,  "3", "syr2k", T,  [S,D,C,Z], ["n","k"], ["layout","triangle","ab_transpose"], ["a","b"], ["c"], ["alpha","beta"], False, "Rank-2K update of a symmetric matrix"),
-  Routine(True,  "3", "her2k", TU, [Ccs,Zzd], ["n","k"], ["layout","triangle","ab_transpose"], ["a","b"], ["c"], ["alpha","beta"], False, "Rank-2K update of a hermitian matrix"),
-  Routine(True,  "3", "trmm",  T,  [S,D,C,Z], ["m","n"], ["layout","side","triangle","a_transpose","diagonal"], ["a"], ["b"], ["alpha"], False, "Triangular matrix-matrix multiplication"),
-  Routine(False, "3", "trsm",  T,  [S,D,C,Z], ["m","n"], ["layout","side","triangle","a_transpose","diagonal"], ["a"], ["b"], ["alpha"], False, "Solves a triangular system of equations"),
+  Routine(True,  "3", "gemm",  T,  [S,D,C,Z], ["m","n","k"], ["layout","a_transpose","b_transpose"], ["a","b"], ["c"], ["alpha","beta"], "", "General matrix-matrix multiplication"),
+  Routine(True,  "3", "symm",  T,  [S,D,C,Z], ["m","n"], ["layout","side","triangle"], ["a","b"], ["c"], ["alpha","beta"], "", "Symmetric matrix-matrix multiplication"),
+  Routine(True,  "3", "hemm",  T,  [C,Z],     ["m","n"], ["layout","side","triangle"], ["a","b"], ["c"], ["alpha","beta"], "", "Hermitian matrix-matrix multiplication"),
+  Routine(True,  "3", "syrk",  T,  [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose"], ["a"], ["c"], ["alpha","beta"], "", "Rank-K update of a symmetric matrix"),
+  Routine(True,  "3", "herk",  Tc, [Css,Zdd], ["n","k"], ["layout","triangle","a_transpose"], ["a"], ["c"], ["alpha","beta"], "", "Rank-K update of a hermitian matrix"),
+  Routine(True,  "3", "syr2k", T,  [S,D,C,Z], ["n","k"], ["layout","triangle","ab_transpose"], ["a","b"], ["c"], ["alpha","beta"], "", "Rank-2K update of a symmetric matrix"),
+  Routine(True,  "3", "her2k", TU, [Ccs,Zzd], ["n","k"], ["layout","triangle","ab_transpose"], ["a","b"], ["c"], ["alpha","beta"], "", "Rank-2K update of a hermitian matrix"),
+  Routine(True,  "3", "trmm",  T,  [S,D,C,Z], ["m","n"], ["layout","side","triangle","a_transpose","diagonal"], ["a"], ["b"], ["alpha"], "", "Triangular matrix-matrix multiplication"),
+  Routine(False, "3", "trsm",  T,  [S,D,C,Z], ["m","n"], ["layout","side","triangle","a_transpose","diagonal"], ["a"], ["b"], ["alpha"], "", "Solves a triangular system of equations"),
 ]]
 
 # ==================================================================================================
@@ -225,7 +228,7 @@ def wrapper_clblas(routines):
 			if routine.scratch:
 				result += "  auto queue = Queue(queues[0]);\n"
 				result += "  auto context = queue.GetContext();\n"
-				result += "  auto scratch_buffer = Buffer<"+flavour.template+">(context, n*x_inc + x_offset);\n"
+				result += "  auto scratch_buffer = Buffer<"+flavour.template+">(context, "+routine.scratch+");\n"
 				arguments += ["scratch_buffer()"]
 			result += "  return clblas"+flavour.name+routine.name+"("
 			result += (",\n"+indent).join([a for a in arguments])
@@ -249,7 +252,7 @@ files = [
   path_clblast+"/src/clblast_c.cc",
   path_clblast+"/test/wrapper_clblas.h",
 ]
-header_lines = [84, 64, 93, 22, 22]
+header_lines = [84, 65, 93, 22, 22]
 footer_lines = [6, 3, 9, 2, 6]
 
 # Checks whether the command-line arguments are valid; exists otherwise
@@ -333,7 +336,7 @@ for level in [1,2,3]:
 				body += "    case clblast::Precision::k"+PrecisionToFullName(precision)+":"
 				found = False
 				for flavour in routine.flavours:
-					if flavour.name == precision:
+					if flavour.precision_name == precision:
 						body += "\n      clblast::RunClient<clblast::TestX"+routine.name+flavour.TestTemplate()
 						body += ">(argc, argv); break;\n"
 						found = True
diff --git a/scripts/generator/routine.py b/scripts/generator/routine.py
index 0a61490b..02040583 100644
--- a/scripts/generator/routine.py
+++ b/scripts/generator/routine.py
@@ -60,7 +60,7 @@ class Routine():
 
 	# List of scalar buffers
 	def ScalarBuffersFirst(self):
-		return ["dot"]
+		return ["dot","nrm2"]
 	def ScalarBuffersSecond(self):
 		return ["sa","sb","sc","ss","sd1","sd2","sx1","sy1","sparam"]
 
@@ -327,7 +327,6 @@ class Routine():
 		        list(chain(*[self.BufferType(b) for b in self.ScalarBuffersSecond()])) +
 		        list(chain(*[self.ScalarType(s, flavour) for s in self.OtherScalars()])))
 
-
 	# ==============================================================================================
 
 	# Retrieves the C++ templated definition for a routine
diff --git a/src/clblast.cc b/src/clblast.cc
index 1b2c3a12..fc50ffae 100644
--- a/src/clblast.cc
+++ b/src/clblast.cc
@@ -26,6 +26,7 @@
 #include "internal/routines/level1/xdot.h"
 #include "internal/routines/level1/xdotu.h"
 #include "internal/routines/level1/xdotc.h"
+#include "internal/routines/level1/xnrm2.h"
 
 // BLAS level-2 includes
 #include "internal/routines/level2/xgemv.h"
@@ -371,6 +372,38 @@ template StatusCode PUBLIC_API Dotc<double2>(const size_t,
                                              const cl_mem, const size_t, const size_t,
                                              cl_command_queue*, cl_event*);
 
+// Euclidian norm of a vector: SNRM2/DNRM2/ScNRM2/DzNRM2
+template <typename T>
+StatusCode Nrm2(const size_t n,
+                cl_mem nrm2_buffer, const size_t nrm2_offset,
+                const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                cl_command_queue* queue, cl_event* event) {
+  auto queue_cpp = Queue(*queue);
+  auto event_cpp = Event(event);
+  auto routine = Xnrm2<T>(queue_cpp, event_cpp);
+  auto status = routine.SetUp();
+  if (status != StatusCode::kSuccess) { return status; }
+  return routine.DoNrm2(n,
+                        Buffer<T>(nrm2_buffer), nrm2_offset,
+                        Buffer<T>(x_buffer), x_offset, x_inc);
+}
+template StatusCode PUBLIC_API Nrm2<float>(const size_t,
+                                           cl_mem, const size_t,
+                                           const cl_mem, const size_t, const size_t,
+                                           cl_command_queue*, cl_event*);
+template StatusCode PUBLIC_API Nrm2<double>(const size_t,
+                                            cl_mem, const size_t,
+                                            const cl_mem, const size_t, const size_t,
+                                            cl_command_queue*, cl_event*);
+template StatusCode PUBLIC_API Nrm2<float2>(const size_t,
+                                            cl_mem, const size_t,
+                                            const cl_mem, const size_t, const size_t,
+                                            cl_command_queue*, cl_event*);
+template StatusCode PUBLIC_API Nrm2<double2>(const size_t,
+                                             cl_mem, const size_t,
+                                             const cl_mem, const size_t, const size_t,
+                                             cl_command_queue*, cl_event*);
+
 // =================================================================================================
 // BLAS level-2 (matrix-vector) routines
 // =================================================================================================
diff --git a/src/clblast_c.cc b/src/clblast_c.cc
index b530732c..6d10c686 100644
--- a/src/clblast_c.cc
+++ b/src/clblast_c.cc
@@ -391,6 +391,48 @@ StatusCode CLBlastZdotc(const size_t n,
   return static_cast<StatusCode>(status);
 }
 
+// NRM2
+StatusCode CLBlastSnrm2(const size_t n,
+                        cl_mem nrm2_buffer, const size_t nrm2_offset,
+                        const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                        cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Nrm2<float>(n,
+                                     nrm2_buffer, nrm2_offset,
+                                     x_buffer, x_offset, x_inc,
+                                     queue, event);
+  return static_cast<StatusCode>(status);
+}
+StatusCode CLBlastDnrm2(const size_t n,
+                        cl_mem nrm2_buffer, const size_t nrm2_offset,
+                        const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                        cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Nrm2<double>(n,
+                                      nrm2_buffer, nrm2_offset,
+                                      x_buffer, x_offset, x_inc,
+                                      queue, event);
+  return static_cast<StatusCode>(status);
+}
+StatusCode CLBlastScnrm2(const size_t n,
+                        cl_mem nrm2_buffer, const size_t nrm2_offset,
+                        const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                        cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Nrm2<float2>(n,
+                                      nrm2_buffer, nrm2_offset,
+                                      x_buffer, x_offset, x_inc,
+                                      queue, event);
+  return static_cast<StatusCode>(status);
+}
+StatusCode CLBlastDznrm2(const size_t n,
+                        cl_mem nrm2_buffer, const size_t nrm2_offset,
+                        const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                        cl_command_queue* queue, cl_event* event) {
+  auto status = clblast::Nrm2<double2>(n,
+                                       nrm2_buffer, nrm2_offset,
+                                       x_buffer, x_offset, x_inc,
+                                       queue, event);
+  return static_cast<StatusCode>(status);
+}
+
 // =================================================================================================
 // BLAS level-2 (matrix-vector) routines
 // =================================================================================================
diff --git a/src/kernels/level1/xnrm2.opencl b/src/kernels/level1/xnrm2.opencl
new file mode 100644
index 00000000..cf579457
--- /dev/null
+++ b/src/kernels/level1/xnrm2.opencl
@@ -0,0 +1,109 @@
+
+// =================================================================================================
+// 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 Xnrm2 kernel. It implements a dot-product computation using reduction
+// kernels. Reduction is split in two parts. In the first (main) kernel the X and Y vectors are
+// multiplied, followed by a per-thread and a per-workgroup reduction. The second (epilogue) kernel
+// is executed with a single workgroup only, computing the final result.
+//
+// =================================================================================================
+
+// Enables loading of this file using the C++ pre-processor's #include (C++11 standard raw string
+// literal). Comment-out this line for syntax-highlighting when developing.
+R"(
+
+// Parameters set by the tuner or by the database. Here they are given a basic default value in case
+// this kernel file is used outside of the CLBlast library.
+#ifndef WGS1
+  #define WGS1 64     // The local work-group size of the main kernel
+#endif
+#ifndef WGS2
+  #define WGS2 64     // The local work-group size of the epilogue kernel
+#endif
+
+// =================================================================================================
+
+// The main reduction kernel, performing the multiplication and the majority of the sum operation
+__attribute__((reqd_work_group_size(WGS1, 1, 1)))
+__kernel void Xnrm2(const int n,
+                    const __global real* restrict xgm, const int x_offset, const int x_inc,
+                    __global real* output) {
+  __local real lm[WGS1];
+  const int lid = get_local_id(0);
+  const int wgid = get_group_id(0);
+  const int num_groups = get_num_groups(0);
+
+  // Performs multiplication and the first steps of the reduction
+  real acc;
+  SetToZero(acc);
+  int id = wgid*WGS1 + lid;
+  while (id < n) {
+    real x1 = xgm[id*x_inc + x_offset];
+    real x2 = x1;
+    COMPLEX_CONJUGATE(x2);
+    MultiplyAdd(acc, x1, x2);
+    id += WGS1*num_groups;
+  }
+  lm[lid] = acc;
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  // Performs reduction in local memory
+  #pragma unroll
+  for (int s=WGS1/2; s>0; s=s>>1) {
+    if (lid < s) {
+      Add(lm[lid], lm[lid], lm[lid + s]);
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+  }
+
+  // Stores the per-workgroup result
+  if (lid == 0) {
+    output[wgid] = lm[0];
+  }
+}
+
+// =================================================================================================
+
+// The epilogue reduction kernel, performing the final bit of the sum operation. This kernel has to
+// be launched with a single workgroup only.
+__attribute__((reqd_work_group_size(WGS2, 1, 1)))
+__kernel void Xnrm2Epilogue(const __global real* restrict input,
+                            __global real* nrm2, const int nrm2_offset) {
+  __local real lm[WGS2];
+  const int lid = get_local_id(0);
+
+  // Performs the first step of the reduction while loading the data
+  Add(lm[lid], input[lid], input[lid + WGS2]);
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  // Performs reduction in local memory
+  #pragma unroll
+  for (int s=WGS2/2; s>0; s=s>>1) {
+    if (lid < s) {
+      Add(lm[lid], lm[lid], lm[lid + s]);
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+  }
+
+  // Computes the square root and stores the final result
+  if (lid == 0) {
+    #if PRECISION == 3232 || PRECISION == 6464
+      nrm2[nrm2_offset].x = sqrt(lm[0].x); // the result is a non-complex number
+    #else
+      nrm2[nrm2_offset] = sqrt(lm[0]);
+    #endif
+  }
+}
+
+// =================================================================================================
+
+// End of the C++11 raw string literal
+)"
+
+// =================================================================================================
diff --git a/src/routines/level1/xnrm2.cc b/src/routines/level1/xnrm2.cc
new file mode 100644
index 00000000..064e68bf
--- /dev/null
+++ b/src/routines/level1/xnrm2.cc
@@ -0,0 +1,107 @@
+
+// =================================================================================================
+// 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 Xnrm2 class (see the header for information about the class).
+//
+// =================================================================================================
+
+#include "internal/routines/level1/xnrm2.h"
+
+#include <string>
+#include <vector>
+
+namespace clblast {
+// =================================================================================================
+
+// Specific implementations to get the memory-type based on a template argument
+template <> const Precision Xnrm2<float>::precision_ = Precision::kSingle;
+template <> const Precision Xnrm2<double>::precision_ = Precision::kDouble;
+template <> const Precision Xnrm2<float2>::precision_ = Precision::kComplexSingle;
+template <> const Precision Xnrm2<double2>::precision_ = Precision::kComplexDouble;
+
+// =================================================================================================
+
+// Constructor: forwards to base class constructor
+template <typename T>
+Xnrm2<T>::Xnrm2(Queue &queue, Event &event, const std::string &name):
+    Routine<T>(queue, event, name, {"Xdot"}, precision_) {
+  source_string_ =
+    #include "../../kernels/level1/xnrm2.opencl"
+  ;
+}
+
+// =================================================================================================
+
+// The main routine
+template <typename T>
+StatusCode Xnrm2<T>::DoNrm2(const size_t n,
+                            const Buffer<T> &nrm2_buffer, const size_t nrm2_offset,
+                            const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc) {
+
+  // Makes sure all dimensions are larger than zero
+  if (n == 0) { return StatusCode::kInvalidDimension; }
+
+  // Tests the vectors for validity
+  auto status = TestVectorX(n, x_buffer, x_offset, x_inc, sizeof(T));
+  if (ErrorIn(status)) { return status; }
+  status = TestVectorDot(1, nrm2_buffer, nrm2_offset, sizeof(T));
+  if (ErrorIn(status)) { return status; }
+
+  // Retrieves the Xnrm2 kernels from the compiled binary
+  try {
+    auto& program = GetProgramFromCache();
+    auto kernel1 = Kernel(program, "Xnrm2");
+    auto kernel2 = Kernel(program, "Xnrm2Epilogue");
+
+    // Creates the buffer for intermediate values
+    auto temp_size = 2*db_["WGS2"];
+    auto temp_buffer = Buffer<T>(context_, temp_size);
+
+    // Sets the kernel arguments
+    kernel1.SetArgument(0, static_cast<int>(n));
+    kernel1.SetArgument(1, x_buffer());
+    kernel1.SetArgument(2, static_cast<int>(x_offset));
+    kernel1.SetArgument(3, static_cast<int>(x_inc));
+    kernel1.SetArgument(4, temp_buffer());
+
+    // Launches the main kernel
+    auto global1 = std::vector<size_t>{db_["WGS1"]*temp_size};
+    auto local1 = std::vector<size_t>{db_["WGS1"]};
+    status = RunKernel(kernel1, global1, local1);
+    if (ErrorIn(status)) { return status; }
+
+    // Sets the arguments for the epilogue kernel
+    kernel2.SetArgument(0, temp_buffer());
+    kernel2.SetArgument(1, nrm2_buffer());
+    kernel2.SetArgument(2, static_cast<int>(nrm2_offset));
+
+    // Launches the epilogue kernel
+    auto global2 = std::vector<size_t>{db_["WGS2"]};
+    auto local2 = std::vector<size_t>{db_["WGS2"]};
+    status = RunKernel(kernel2, global2, local2);
+    if (ErrorIn(status)) { return status; }
+
+    // Waits for all kernels to finish
+    queue_.Finish();
+
+    // Succesfully finished the computation
+    return StatusCode::kSuccess;
+  } catch (...) { return StatusCode::kInvalidKernel; }
+}
+
+// =================================================================================================
+
+// Compiles the templated class
+template class Xnrm2<float>;
+template class Xnrm2<double>;
+template class Xnrm2<float2>;
+template class Xnrm2<double2>;
+
+// =================================================================================================
+} // namespace clblast
diff --git a/test/correctness/routines/level1/xnrm2.cc b/test/correctness/routines/level1/xnrm2.cc
new file mode 100644
index 00000000..97fb0ad6
--- /dev/null
+++ b/test/correctness/routines/level1/xnrm2.cc
@@ -0,0 +1,28 @@
+
+// =================================================================================================
+// 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>
+//
+// =================================================================================================
+
+#include "correctness/testblas.h"
+#include "routines/level1/xnrm2.h"
+
+// 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[]) {
+  clblast::RunTests<clblast::TestXnrm2<float>, float, float>(argc, argv, false, "SNRM2");
+  clblast::RunTests<clblast::TestXnrm2<double>, double, double>(argc, argv, true, "DNRM2");
+  clblast::RunTests<clblast::TestXnrm2<float2>, float2, float2>(argc, argv, true, "ScNRM2");
+  clblast::RunTests<clblast::TestXnrm2<double2>, double2, double2>(argc, argv, true, "DzNRM2");
+  return 0;
+}
+
+// =================================================================================================
diff --git a/test/correctness/testblas.cc b/test/correctness/testblas.cc
index febd7504..1329b2c5 100644
--- a/test/correctness/testblas.cc
+++ b/test/correctness/testblas.cc
@@ -58,14 +58,14 @@ TestBlas<T,U>::TestBlas(int argc, char *argv[], const bool silent,
   b_source_.resize(std::max(max_mat, max_matvec)*std::max(max_ld, max_matvec) + max_offset);
   c_source_.resize(std::max(max_mat, max_matvec)*std::max(max_ld, max_matvec) + max_offset);
   ap_source_.resize(std::max(max_mat, max_matvec)*std::max(max_mat, max_matvec) + max_offset);
-  dot_source_.resize(std::max(max_mat, max_matvec) + max_offset);
+  scalar_source_.resize(std::max(max_mat, max_matvec) + max_offset);
   PopulateVector(x_source_);
   PopulateVector(y_source_);
   PopulateVector(a_source_);
   PopulateVector(b_source_);
   PopulateVector(c_source_);
   PopulateVector(ap_source_);
-  PopulateVector(dot_source_);
+  PopulateVector(scalar_source_);
 }
 
 // ===============================================================================================
@@ -86,15 +86,15 @@ void TestBlas<T,U>::TestRegular(std::vector<Arguments<U>> &test_vector, const st
     auto b_mat1 = Buffer<T>(context_, args.b_size);
     auto c_mat1 = Buffer<T>(context_, args.c_size);
     auto ap_mat1 = Buffer<T>(context_, args.ap_size);
-    auto dot1 = Buffer<T>(context_, args.dot_size);
+    auto scalar1 = Buffer<T>(context_, args.scalar_size);
     x_vec1.Write(queue_, args.x_size, x_source_);
     y_vec1.Write(queue_, args.y_size, y_source_);
     a_mat1.Write(queue_, args.a_size, a_source_);
     b_mat1.Write(queue_, args.b_size, b_source_);
     c_mat1.Write(queue_, args.c_size, c_source_);
     ap_mat1.Write(queue_, args.ap_size, ap_source_);
-    dot1.Write(queue_, args.dot_size, dot_source_);
-    auto buffers1 = Buffers<T>{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, ap_mat1, dot1};
+    scalar1.Write(queue_, args.scalar_size, scalar_source_);
+    auto buffers1 = Buffers<T>{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, ap_mat1, scalar1};
     auto status1 = run_reference_(args, buffers1, queue_);
 
     // Runs the CLBlast code
@@ -104,15 +104,15 @@ void TestBlas<T,U>::TestRegular(std::vector<Arguments<U>> &test_vector, const st
     auto b_mat2 = Buffer<T>(context_, args.b_size);
     auto c_mat2 = Buffer<T>(context_, args.c_size);
     auto ap_mat2 = Buffer<T>(context_, args.ap_size);
-    auto dot2 = Buffer<T>(context_, args.dot_size);
+    auto scalar2 = Buffer<T>(context_, args.scalar_size);
     x_vec2.Write(queue_, args.x_size, x_source_);
     y_vec2.Write(queue_, args.y_size, y_source_);
     a_mat2.Write(queue_, args.a_size, a_source_);
     b_mat2.Write(queue_, args.b_size, b_source_);
     c_mat2.Write(queue_, args.c_size, c_source_);
     ap_mat2.Write(queue_, args.ap_size, ap_source_);
-    dot2.Write(queue_, args.dot_size, dot_source_);
-    auto buffers2 = Buffers<T>{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, ap_mat2, dot2};
+    scalar2.Write(queue_, args.scalar_size, scalar_source_);
+    auto buffers2 = Buffers<T>{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, ap_mat2, scalar2};
     auto status2 = run_routine_(args, buffers2, queue_);
 
     // Tests for equality of the two status codes
@@ -162,32 +162,32 @@ void TestBlas<T,U>::TestInvalid(std::vector<Arguments<U>> &test_vector, const st
     auto b1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.b_size*sizeof(T), nullptr,nullptr);
     auto c1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.c_size*sizeof(T), nullptr,nullptr);
     auto ap1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.ap_size*sizeof(T), nullptr,nullptr);
-    auto d1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.dot_size*sizeof(T), nullptr,nullptr);
+    auto d1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.scalar_size*sizeof(T), nullptr,nullptr);
     auto x_vec1 = Buffer<T>(x1);
     auto y_vec1 = Buffer<T>(y1);
     auto a_mat1 = Buffer<T>(a1);
     auto b_mat1 = Buffer<T>(b1);
     auto c_mat1 = Buffer<T>(c1);
     auto ap_mat1 = Buffer<T>(ap1);
-    auto dot1 = Buffer<T>(d1);
+    auto scalar1 = Buffer<T>(d1);
     auto x2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.x_size*sizeof(T), nullptr,nullptr);
     auto y2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.y_size*sizeof(T), nullptr,nullptr);
     auto a2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.a_size*sizeof(T), nullptr,nullptr);
     auto b2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.b_size*sizeof(T), nullptr,nullptr);
     auto c2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.c_size*sizeof(T), nullptr,nullptr);
     auto ap2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.ap_size*sizeof(T), nullptr,nullptr);
-    auto d2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.dot_size*sizeof(T), nullptr,nullptr);
+    auto d2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.scalar_size*sizeof(T), nullptr,nullptr);
     auto x_vec2 = Buffer<T>(x2);
     auto y_vec2 = Buffer<T>(y2);
     auto a_mat2 = Buffer<T>(a2);
     auto b_mat2 = Buffer<T>(b2);
     auto c_mat2 = Buffer<T>(c2);
     auto ap_mat2 = Buffer<T>(ap2);
-    auto dot2 = Buffer<T>(d2);
+    auto scalar2 = Buffer<T>(d2);
 
     // Runs the two routines
-    auto buffers1 = Buffers<T>{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, ap_mat1, dot1};
-    auto buffers2 = Buffers<T>{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, ap_mat2, dot2};
+    auto buffers1 = Buffers<T>{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, ap_mat1, scalar1};
+    auto buffers2 = Buffers<T>{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, ap_mat2, scalar2};
     auto status1 = run_reference_(args, buffers1, queue_);
     auto status2 = run_routine_(args, buffers2, queue_);
 
diff --git a/test/correctness/testblas.h b/test/correctness/testblas.h
index bfd1763c..7c9032bd 100644
--- a/test/correctness/testblas.h
+++ b/test/correctness/testblas.h
@@ -92,7 +92,7 @@ class TestBlas: public Tester<T,U> {
   std::vector<T> b_source_;
   std::vector<T> c_source_;
   std::vector<T> ap_source_;
-  std::vector<T> dot_source_;
+  std::vector<T> scalar_source_;
   
   // The routine-specific functions passed to the tester
   Routine run_routine_;
@@ -143,6 +143,7 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name
   auto c_offsets = std::vector<size_t>{args.c_offset};
   auto ap_offsets = std::vector<size_t>{args.ap_offset};
   auto dot_offsets = std::vector<size_t>{args.dot_offset};
+  auto nrm2_offsets = std::vector<size_t>{args.nrm2_offset};
   auto alphas = std::vector<U>{args.alpha};
   auto betas = std::vector<U>{args.beta};
   auto x_sizes = std::vector<size_t>{args.x_size};
@@ -182,6 +183,7 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name
     if (option == kArgCOffset) { c_offsets = tester.kOffsets; }
     if (option == kArgAPOffset) { ap_offsets = tester.kOffsets; }
     if (option == kArgDotOffset) { dot_offsets = tester.kOffsets; }
+    if (option == kArgNrm2Offset) { nrm2_offsets = tester.kOffsets; }
     if (option == kArgAlpha) { alphas = tester.kAlphaValues; }
     if (option == kArgBeta) { betas = tester.kBetaValues; }
 
@@ -221,10 +223,12 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name
                                           for (auto &c_offset: c_offsets) { r_args.c_offset = c_offset;
                                             for (auto &ap_offset: ap_offsets) { r_args.ap_offset = ap_offset;
                                               for (auto &dot_offset: dot_offsets) { r_args.dot_offset = dot_offset;
-                                                for (auto &alpha: alphas) { r_args.alpha = alpha;
-                                                  for (auto &beta: betas) { r_args.beta = beta;
-                                                    C::SetSizes(r_args);
-                                                    regular_test_vector.push_back(r_args);
+                                                for (auto &nrm2_offset: nrm2_offsets) { r_args.nrm2_offset = nrm2_offset;
+                                                  for (auto &alpha: alphas) { r_args.alpha = alpha;
+                                                    for (auto &beta: betas) { r_args.beta = beta;
+                                                      C::SetSizes(r_args);
+                                                      regular_test_vector.push_back(r_args);
+                                                    }
                                                   }
                                                 }
                                               }
diff --git a/test/performance/client.cc b/test/performance/client.cc
index ebfad3a6..17f54231 100644
--- a/test/performance/client.cc
+++ b/test/performance/client.cc
@@ -136,14 +136,14 @@ void Client<T,U>::PerformanceTest(Arguments<U> &args, const SetMetric set_sizes)
     std::vector<T> b_source(args.b_size);
     std::vector<T> c_source(args.c_size);
     std::vector<T> ap_source(args.ap_size);
-    std::vector<T> dot_source(args.dot_size);
+    std::vector<T> scalar_source(args.scalar_size);
     PopulateVector(x_source);
     PopulateVector(y_source);
     PopulateVector(a_source);
     PopulateVector(b_source);
     PopulateVector(c_source);
     PopulateVector(ap_source);
-    PopulateVector(dot_source);
+    PopulateVector(scalar_source);
 
     // Creates the matrices on the device
     auto x_vec = Buffer<T>(context, args.x_size);
@@ -152,15 +152,15 @@ void Client<T,U>::PerformanceTest(Arguments<U> &args, const SetMetric set_sizes)
     auto b_mat = Buffer<T>(context, args.b_size);
     auto c_mat = Buffer<T>(context, args.c_size);
     auto ap_mat = Buffer<T>(context, args.ap_size);
-    auto dot = Buffer<T>(context, args.dot_size);
+    auto scalar = Buffer<T>(context, args.scalar_size);
     x_vec.Write(queue, args.x_size, x_source);
     y_vec.Write(queue, args.y_size, y_source);
     a_mat.Write(queue, args.a_size, a_source);
     b_mat.Write(queue, args.b_size, b_source);
     c_mat.Write(queue, args.c_size, c_source);
     ap_mat.Write(queue, args.ap_size, ap_source);
-    dot.Write(queue, args.dot_size, dot_source);
-    auto buffers = Buffers<T>{x_vec, y_vec, a_mat, b_mat, c_mat, ap_mat, dot};
+    scalar.Write(queue, args.scalar_size, scalar_source);
+    auto buffers = Buffers<T>{x_vec, y_vec, a_mat, b_mat, c_mat, ap_mat, scalar};
 
     // Runs the routines and collects the timings
     auto timings = std::vector<std::pair<std::string, double>>();
@@ -267,6 +267,7 @@ void Client<T,U>::PrintTableRow(const Arguments<U>& args,
     else if (o == kArgCOffset) {  integers.push_back(args.c_offset); }
     else if (o == kArgAPOffset) { integers.push_back(args.ap_offset); }
     else if (o == kArgDotOffset) {integers.push_back(args.dot_offset); }
+    else if (o == kArgNrm2Offset){integers.push_back(args.nrm2_offset); }
   }
   auto strings = std::vector<std::string>{};
   for (auto &o: options_) {
diff --git a/test/performance/routines/level1/xnrm2.cc b/test/performance/routines/level1/xnrm2.cc
new file mode 100644
index 00000000..db6ec9ad
--- /dev/null
+++ b/test/performance/routines/level1/xnrm2.cc
@@ -0,0 +1,35 @@
+
+// =================================================================================================
+// 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>
+//
+// =================================================================================================
+
+#include "performance/client.h"
+#include "routines/level1/xnrm2.h"
+
+// 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, clblast::Precision::kSingle)) {
+    case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode");
+    case clblast::Precision::kSingle:
+      clblast::RunClient<clblast::TestXnrm2<float>, float, float>(argc, argv); break;
+    case clblast::Precision::kDouble:
+      clblast::RunClient<clblast::TestXnrm2<double>, double, double>(argc, argv); break;
+    case clblast::Precision::kComplexSingle:
+      clblast::RunClient<clblast::TestXnrm2<float2>, float2, float2>(argc, argv); break;
+    case clblast::Precision::kComplexDouble:
+      clblast::RunClient<clblast::TestXnrm2<double2>, double2, double2>(argc, argv); break;
+  }
+  return 0;
+}
+
+// =================================================================================================
diff --git a/test/routines/level1/xdot.h b/test/routines/level1/xdot.h
index bfcfdaff..04669f52 100644
--- a/test/routines/level1/xdot.h
+++ b/test/routines/level1/xdot.h
@@ -54,7 +54,7 @@ class TestXdot {
   static void SetSizes(Arguments<T> &args) {
     args.x_size = GetSizeX(args);
     args.y_size = GetSizeY(args);
-    args.dot_size = GetSizeDot(args);
+    args.scalar_size = GetSizeDot(args);
   }
 
   // Describes what the default values of the leading dimensions of the matrices are
@@ -72,7 +72,7 @@ class TestXdot {
     auto queue_plain = queue();
     auto event = cl_event{};
     auto status = Dot<T>(args.n,
-                         buffers.dot(), args.dot_offset,
+                         buffers.scalar(), args.dot_offset,
                          buffers.x_vec(), args.x_offset, args.x_inc,
                          buffers.y_vec(), args.y_offset, args.y_inc,
                          &queue_plain, &event);
@@ -85,7 +85,7 @@ class TestXdot {
     auto queue_plain = queue();
     auto event = cl_event{};
     auto status = clblasXdot<T>(args.n,
-                                buffers.dot(), args.dot_offset,
+                                buffers.scalar(), args.dot_offset,
                                 buffers.x_vec(), args.x_offset, args.x_inc,
                                 buffers.y_vec(), args.y_offset, args.y_inc,
                                 1, &queue_plain, 0, nullptr, &event);
@@ -95,8 +95,8 @@ class TestXdot {
 
   // 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.dot_size, static_cast<T>(0));
-    buffers.dot.Read(queue, args.dot_size, result);
+    std::vector<T> result(args.scalar_size, static_cast<T>(0));
+    buffers.scalar.Read(queue, args.scalar_size, result);
     return result;
   }
 
diff --git a/test/routines/level1/xdotc.h b/test/routines/level1/xdotc.h
index e403ba4c..e5b42ef4 100644
--- a/test/routines/level1/xdotc.h
+++ b/test/routines/level1/xdotc.h
@@ -54,7 +54,7 @@ class TestXdotc {
   static void SetSizes(Arguments<T> &args) {
     args.x_size = GetSizeX(args);
     args.y_size = GetSizeY(args);
-    args.dot_size = GetSizeDot(args);
+    args.scalar_size = GetSizeDot(args);
   }
 
   // Describes what the default values of the leading dimensions of the matrices are
@@ -72,7 +72,7 @@ class TestXdotc {
     auto queue_plain = queue();
     auto event = cl_event{};
     auto status = Dotc<T>(args.n,
-                          buffers.dot(), args.dot_offset,
+                          buffers.scalar(), args.dot_offset,
                           buffers.x_vec(), args.x_offset, args.x_inc,
                           buffers.y_vec(), args.y_offset, args.y_inc,
                           &queue_plain, &event);
@@ -85,7 +85,7 @@ class TestXdotc {
     auto queue_plain = queue();
     auto event = cl_event{};
     auto status = clblasXdotc<T>(args.n,
-                                 buffers.dot(), args.dot_offset,
+                                 buffers.scalar(), args.dot_offset,
                                  buffers.x_vec(), args.x_offset, args.x_inc,
                                  buffers.y_vec(), args.y_offset, args.y_inc,
                                  1, &queue_plain, 0, nullptr, &event);
@@ -95,8 +95,8 @@ class TestXdotc {
 
   // 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.dot_size, static_cast<T>(0));
-    buffers.dot.Read(queue, args.dot_size, result);
+    std::vector<T> result(args.scalar_size, static_cast<T>(0));
+    buffers.scalar.Read(queue, args.scalar_size, result);
     return result;
   }
 
diff --git a/test/routines/level1/xdotu.h b/test/routines/level1/xdotu.h
index 8b2c65a8..6430148c 100644
--- a/test/routines/level1/xdotu.h
+++ b/test/routines/level1/xdotu.h
@@ -54,7 +54,7 @@ class TestXdotu {
   static void SetSizes(Arguments<T> &args) {
     args.x_size = GetSizeX(args);
     args.y_size = GetSizeY(args);
-    args.dot_size = GetSizeDot(args);
+    args.scalar_size = GetSizeDot(args);
   }
 
   // Describes what the default values of the leading dimensions of the matrices are
@@ -72,7 +72,7 @@ class TestXdotu {
     auto queue_plain = queue();
     auto event = cl_event{};
     auto status = Dotu<T>(args.n,
-                          buffers.dot(), args.dot_offset,
+                          buffers.scalar(), args.dot_offset,
                           buffers.x_vec(), args.x_offset, args.x_inc,
                           buffers.y_vec(), args.y_offset, args.y_inc,
                           &queue_plain, &event);
@@ -85,7 +85,7 @@ class TestXdotu {
     auto queue_plain = queue();
     auto event = cl_event{};
     auto status = clblasXdotu<T>(args.n,
-                                 buffers.dot(), args.dot_offset,
+                                 buffers.scalar(), args.dot_offset,
                                  buffers.x_vec(), args.x_offset, args.x_inc,
                                  buffers.y_vec(), args.y_offset, args.y_inc,
                                  1, &queue_plain, 0, nullptr, &event);
@@ -95,8 +95,8 @@ class TestXdotu {
 
   // 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.dot_size, static_cast<T>(0));
-    buffers.dot.Read(queue, args.dot_size, result);
+    std::vector<T> result(args.scalar_size, static_cast<T>(0));
+    buffers.scalar.Read(queue, args.scalar_size, result);
     return result;
   }
 
diff --git a/test/routines/level1/xnrm2.h b/test/routines/level1/xnrm2.h
new file mode 100644
index 00000000..e3f77ee4
--- /dev/null
+++ b/test/routines/level1/xnrm2.h
@@ -0,0 +1,117 @@
+
+// =================================================================================================
+// 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 Xnrm2 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_XNRM2_H_
+#define CLBLAST_TEST_ROUTINES_XNRM2_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 TestXnrm2 {
+ public:
+
+  // The BLAS level: 1, 2, or 3
+  static size_t BLASLevel() { return 1; }
+
+  // The list of arguments relevant for this routine
+  static std::vector<std::string> GetOptions() {
+    return {kArgN,
+            kArgXInc,
+            kArgXOffset, kArgNrm2Offset};
+  }
+
+  // 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 GetSizeNrm2(const Arguments<T> &args) {
+    return 1 + args.nrm2_offset;
+  }
+
+  // Describes how to set the sizes of all the buffers
+  static void SetSizes(Arguments<T> &args) {
+    args.x_size = GetSizeX(args);
+    args.scalar_size = GetSizeNrm2(args);
+  }
+
+  // Describes what the default values of the leading dimensions of the matrices are
+  static size_t DefaultLDA(const Arguments<T> &) { return 1; } // N/A for this routine
+  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 = Nrm2<T>(args.n,
+                          buffers.scalar(), args.nrm2_offset,
+                          buffers.x_vec(), args.x_offset, args.x_inc,
+                          &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 = clblasXnrm2<T>(args.n,
+                                 buffers.scalar(), args.nrm2_offset,
+                                 buffers.x_vec(), args.x_offset, args.x_inc,
+                                 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.scalar_size, static_cast<T>(0));
+    buffers.scalar.Read(queue, args.scalar_size, result);
+    return result;
+  }
+
+  // Describes how to compute the indices of the result buffer
+  static size_t ResultID1(const Arguments<T> &) { return 1; } // N/A for this routine
+  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, const size_t) {
+    return args.nrm2_offset;
+  }
+
+  // Describes how to compute performance metrics
+  static size_t GetFlops(const Arguments<T> &args) {
+    return 2 * args.n;
+  }
+  static size_t GetBytes(const Arguments<T> &args) {
+    return ((args.n) + 1) * sizeof(T);
+  }
+};
+
+// =================================================================================================
+} // namespace clblast
+
+// CLBLAST_TEST_ROUTINES_XNRM2_H_
+#endif
diff --git a/test/wrapper_clblas.h b/test/wrapper_clblas.h
index 259aa27c..fb6e83aa 100644
--- a/test/wrapper_clblas.h
+++ b/test/wrapper_clblas.h
@@ -378,7 +378,7 @@ clblasStatus clblasXdot<float>(const size_t n,
                                cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float>(context, n);
   return clblasSdot(n,
                     dot_buffer, dot_offset,
                     x_buffer, x_offset, static_cast<int>(x_inc),
@@ -395,7 +395,7 @@ clblasStatus clblasXdot<double>(const size_t n,
                                 cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double>(context, n);
   return clblasDdot(n,
                     dot_buffer, dot_offset,
                     x_buffer, x_offset, static_cast<int>(x_inc),
@@ -421,7 +421,7 @@ clblasStatus clblasXdotu<float2>(const size_t n,
                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float2>(context, n);
   return clblasCdotu(n,
                      dot_buffer, dot_offset,
                      x_buffer, x_offset, static_cast<int>(x_inc),
@@ -438,7 +438,7 @@ clblasStatus clblasXdotu<double2>(const size_t n,
                                   cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double2>(context, n);
   return clblasZdotu(n,
                      dot_buffer, dot_offset,
                      x_buffer, x_offset, static_cast<int>(x_inc),
@@ -464,7 +464,7 @@ clblasStatus clblasXdotc<float2>(const size_t n,
                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float2>(context, n);
   return clblasCdotc(n,
                      dot_buffer, dot_offset,
                      x_buffer, x_offset, static_cast<int>(x_inc),
@@ -481,7 +481,7 @@ clblasStatus clblasXdotc<double2>(const size_t n,
                                   cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double2>(context, n);
   return clblasZdotc(n,
                      dot_buffer, dot_offset,
                      x_buffer, x_offset, static_cast<int>(x_inc),
@@ -490,6 +490,74 @@ clblasStatus clblasXdotc<double2>(const size_t n,
                      num_queues, queues, num_wait_events, wait_events, events);
 }
 
+// Forwards the clBLAS calls for SNRM2/DNRM2/ScNRM2/DzNRM2
+template <typename T>
+clblasStatus clblasXnrm2(const size_t n,
+                         cl_mem nrm2_buffer, const size_t nrm2_offset,
+                         const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                         cl_uint num_queues, cl_command_queue *queues,
+                         cl_uint num_wait_events, const cl_event *wait_events, cl_event *events);
+template <>
+clblasStatus clblasXnrm2<float>(const size_t n,
+                                cl_mem nrm2_buffer, const size_t nrm2_offset,
+                                const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                                cl_uint num_queues, cl_command_queue *queues,
+                                cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
+  auto queue = Queue(queues[0]);
+  auto context = queue.GetContext();
+  auto scratch_buffer = Buffer<float>(context, 2*n);
+  return clblasSnrm2(n,
+                     nrm2_buffer, nrm2_offset,
+                     x_buffer, x_offset, static_cast<int>(x_inc),
+                     scratch_buffer(),
+                     num_queues, queues, num_wait_events, wait_events, events);
+}
+template <>
+clblasStatus clblasXnrm2<double>(const size_t n,
+                                 cl_mem nrm2_buffer, const size_t nrm2_offset,
+                                 const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                                 cl_uint num_queues, cl_command_queue *queues,
+                                 cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
+  auto queue = Queue(queues[0]);
+  auto context = queue.GetContext();
+  auto scratch_buffer = Buffer<double>(context, 2*n);
+  return clblasDnrm2(n,
+                     nrm2_buffer, nrm2_offset,
+                     x_buffer, x_offset, static_cast<int>(x_inc),
+                     scratch_buffer(),
+                     num_queues, queues, num_wait_events, wait_events, events);
+}
+template <>
+clblasStatus clblasXnrm2<float2>(const size_t n,
+                                 cl_mem nrm2_buffer, const size_t nrm2_offset,
+                                 const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                                 cl_uint num_queues, cl_command_queue *queues,
+                                 cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
+  auto queue = Queue(queues[0]);
+  auto context = queue.GetContext();
+  auto scratch_buffer = Buffer<float2>(context, 2*n);
+  return clblasScnrm2(n,
+                     nrm2_buffer, nrm2_offset,
+                     x_buffer, x_offset, static_cast<int>(x_inc),
+                     scratch_buffer(),
+                     num_queues, queues, num_wait_events, wait_events, events);
+}
+template <>
+clblasStatus clblasXnrm2<double2>(const size_t n,
+                                  cl_mem nrm2_buffer, const size_t nrm2_offset,
+                                  const cl_mem x_buffer, const size_t x_offset, const size_t x_inc,
+                                  cl_uint num_queues, cl_command_queue *queues,
+                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
+  auto queue = Queue(queues[0]);
+  auto context = queue.GetContext();
+  auto scratch_buffer = Buffer<double2>(context, 2*n);
+  return clblasDznrm2(n,
+                     nrm2_buffer, nrm2_offset,
+                     x_buffer, x_offset, static_cast<int>(x_inc),
+                     scratch_buffer(),
+                     num_queues, queues, num_wait_events, wait_events, events);
+}
+
 // =================================================================================================
 // BLAS level-2 (matrix-vector) routines
 // =================================================================================================
@@ -887,7 +955,7 @@ clblasStatus clblasXtrmv<float>(const clblasOrder layout, const clblasUplo trian
                                 cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float>(context, n);
   return clblasStrmv(layout, triangle, a_transpose, diagonal,
                      n,
                      a_buffer, a_offset, a_ld,
@@ -904,7 +972,7 @@ clblasStatus clblasXtrmv<double>(const clblasOrder layout, const clblasUplo tria
                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double>(context, n);
   return clblasDtrmv(layout, triangle, a_transpose, diagonal,
                      n,
                      a_buffer, a_offset, a_ld,
@@ -921,7 +989,7 @@ clblasStatus clblasXtrmv<float2>(const clblasOrder layout, const clblasUplo tria
                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float2>(context, n);
   return clblasCtrmv(layout, triangle, a_transpose, diagonal,
                      n,
                      a_buffer, a_offset, a_ld,
@@ -938,7 +1006,7 @@ clblasStatus clblasXtrmv<double2>(const clblasOrder layout, const clblasUplo tri
                                   cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double2>(context, n);
   return clblasZtrmv(layout, triangle, a_transpose, diagonal,
                      n,
                      a_buffer, a_offset, a_ld,
@@ -964,7 +1032,7 @@ clblasStatus clblasXtbmv<float>(const clblasOrder layout, const clblasUplo trian
                                 cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float>(context, n);
   return clblasStbmv(layout, triangle, a_transpose, diagonal,
                      n, k,
                      a_buffer, a_offset, a_ld,
@@ -981,7 +1049,7 @@ clblasStatus clblasXtbmv<double>(const clblasOrder layout, const clblasUplo tria
                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double>(context, n);
   return clblasDtbmv(layout, triangle, a_transpose, diagonal,
                      n, k,
                      a_buffer, a_offset, a_ld,
@@ -998,7 +1066,7 @@ clblasStatus clblasXtbmv<float2>(const clblasOrder layout, const clblasUplo tria
                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float2>(context, n);
   return clblasCtbmv(layout, triangle, a_transpose, diagonal,
                      n, k,
                      a_buffer, a_offset, a_ld,
@@ -1015,7 +1083,7 @@ clblasStatus clblasXtbmv<double2>(const clblasOrder layout, const clblasUplo tri
                                   cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double2>(context, n);
   return clblasZtbmv(layout, triangle, a_transpose, diagonal,
                      n, k,
                      a_buffer, a_offset, a_ld,
@@ -1041,7 +1109,7 @@ clblasStatus clblasXtpmv<float>(const clblasOrder layout, const clblasUplo trian
                                 cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float>(context, n);
   return clblasStpmv(layout, triangle, a_transpose, diagonal,
                      n,
                      ap_buffer, ap_offset,
@@ -1058,7 +1126,7 @@ clblasStatus clblasXtpmv<double>(const clblasOrder layout, const clblasUplo tria
                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double>(context, n);
   return clblasDtpmv(layout, triangle, a_transpose, diagonal,
                      n,
                      ap_buffer, ap_offset,
@@ -1075,7 +1143,7 @@ clblasStatus clblasXtpmv<float2>(const clblasOrder layout, const clblasUplo tria
                                  cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<float2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<float2>(context, n);
   return clblasCtpmv(layout, triangle, a_transpose, diagonal,
                      n,
                      ap_buffer, ap_offset,
@@ -1092,7 +1160,7 @@ clblasStatus clblasXtpmv<double2>(const clblasOrder layout, const clblasUplo tri
                                   cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) {
   auto queue = Queue(queues[0]);
   auto context = queue.GetContext();
-  auto scratch_buffer = Buffer<double2>(context, n*x_inc + x_offset);
+  auto scratch_buffer = Buffer<double2>(context, n);
   return clblasZtpmv(layout, triangle, a_transpose, diagonal,
                      n,
                      ap_buffer, ap_offset,