summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
Diffstat (limited to 'src')
-rw-r--r--src/clblast.cc177
-rw-r--r--src/database.cc2
-rw-r--r--src/kernels/common.opencl7
-rw-r--r--src/kernels/level2/level2.opencl102
-rw-r--r--src/kernels/level2/xger.opencl106
-rw-r--r--src/kernels/level2/xher.opencl73
-rw-r--r--src/routines/level1/xdotu.cc1
-rw-r--r--src/routines/level2/xger.cc112
-rw-r--r--src/routines/level2/xgerc.cc53
-rw-r--r--src/routines/level2/xgeru.cc52
-rw-r--r--src/routines/level2/xher.cc122
-rw-r--r--src/routines/level2/xhpr.cc50
-rw-r--r--src/routines/level2/xspr.cc50
-rw-r--r--src/routines/level2/xsyr.cc50
-rw-r--r--src/tuning/xger.cc129
-rw-r--r--src/utilities.cc4
16 files changed, 1035 insertions, 55 deletions
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);
}
// =================================================================================================
generated by cgit v1.2.3 (git 2.39.1) at 2024-08-03 16:42:32 +0000