diff options
author | CNugteren <web@cedricnugteren.nl> | 2015-06-13 11:01:20 +0200 |
---|---|---|
committer | CNugteren <web@cedricnugteren.nl> | 2015-06-13 11:01:20 +0200 |
commit | e522d1a74e6a877f32730da4807f54cf9a996679 (patch) | |
tree | 68ebcd40b52b3a671b5223885736a569606d5656 | |
parent | 7610a8b182b1579857b134b334228d8cff41f84f (diff) |
Added initial version of GEMV including tester and performance client
-rw-r--r-- | CMakeLists.txt | 2 | ||||
-rw-r--r-- | include/clblast.h | 17 | ||||
-rw-r--r-- | include/internal/database.h | 1 | ||||
-rw-r--r-- | include/internal/database/xgemv.h | 129 | ||||
-rw-r--r-- | include/internal/routines/xgemv.h | 46 | ||||
-rw-r--r-- | src/clblast.cc | 155 | ||||
-rw-r--r-- | src/database.cc | 2 | ||||
-rw-r--r-- | src/kernels/xgemv.opencl | 14 | ||||
-rw-r--r-- | src/routines/xgemv.cc | 117 | ||||
-rw-r--r-- | src/tuning/xgemv.cc | 3 | ||||
-rw-r--r-- | test/correctness/routines/xgemv.cc | 94 | ||||
-rw-r--r-- | test/correctness/testaxy.cc | 26 | ||||
-rw-r--r-- | test/performance/client.cc | 94 | ||||
-rw-r--r-- | test/performance/client.h | 6 | ||||
-rw-r--r-- | test/performance/routines/xgemv.cc | 107 | ||||
-rw-r--r-- | test/wrapper_clblas.h | 58 |
16 files changed, 783 insertions, 88 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index b6697ac3..6a597f22 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -96,7 +96,7 @@ include_directories(${clblast_SOURCE_DIR}/include ${OPENCL_INCLUDE_DIRS}) set(KERNELS copy pad transpose padtranspose xaxpy xgemv xgemm) set(SAMPLE_PROGRAMS sgemm) set(ROUTINES_XY xaxpy) -set(ROUTINES_AXY ) +set(ROUTINES_AXY xgemv) set(ROUTINES_ABC xgemm xsymm) set(ROUTINES ${ROUTINES_XY} ${ROUTINES_AXY} ${ROUTINES_ABC}) diff --git a/include/clblast.h b/include/clblast.h index 4c3c5201..231348b8 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -85,7 +85,7 @@ enum class Precision { kHalf = 16, kSingle = 32, kDouble = 64, // Templated-precision vector-times-constant plus vector: SAXPY/DAXPY/CAXPY/ZAXPY template <typename T> -StatusCode Axpy(const size_t m, const T alpha, +StatusCode Axpy(const size_t n, const T alpha, const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, cl_mem y_buffer, const size_t y_offset, const size_t y_inc, cl_command_queue* queue, cl_event* event); @@ -93,10 +93,21 @@ StatusCode Axpy(const size_t m, const T alpha, // ================================================================================================= // BLAS level-2 (matrix-vector) routines +// Templated-precision generalized matrix-vector multiplication: SGEMV/DGEMV/CGEMV/ZGEMV +template <typename T> +StatusCode Gemv(const Layout layout, const Transpose transpose_a, + const size_t m, const size_t n, + const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const T beta, + cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); + // ================================================================================================= // BLAS level-3 (matrix-matrix) routines -// Templated-precision generalized matrix multiplication: SGEMM/DGEMM +// Templated-precision generalized matrix-matrix multiplication: SGEMM/DGEMM template <typename T> StatusCode Gemm(const Layout layout, const Transpose transpose_a, const Transpose transpose_b, const size_t m, const size_t n, const size_t k, @@ -107,7 +118,7 @@ StatusCode Gemm(const Layout layout, const Transpose transpose_a, const Transpos cl_mem c_buffer, const size_t c_offset, const size_t c_ld, cl_command_queue* queue, cl_event* event); -// Templated-precision symmetric matrix multiplication: SSYMM/DSYMM +// Templated-precision symmetric matrix-matrix multiplication: SSYMM/DSYMM template <typename T> StatusCode Symm(const Layout layout, const Side side, const Triangle triangle, const size_t m, const size_t n, diff --git a/include/internal/database.h b/include/internal/database.h index dbbdd5c0..33ad1979 100644 --- a/include/internal/database.h +++ b/include/internal/database.h @@ -54,6 +54,7 @@ class Database { // The database consists of separate database entries, stored together in a vector static const DatabaseEntry XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble; + static const DatabaseEntry XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble; static const DatabaseEntry XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble; static const DatabaseEntry CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble; static const DatabaseEntry PadSingle, PadDouble, PadComplexSingle, PadComplexDouble; diff --git a/include/internal/database/xgemv.h b/include/internal/database/xgemv.h new file mode 100644 index 00000000..178d1122 --- /dev/null +++ b/include/internal/database/xgemv.h @@ -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 populates the database with best-found tuning parameters for the Xgemv kernels. +// +// ================================================================================================= + +namespace clblast { +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvSingle = { + "Xgemv", Precision::kSingle, { + { // NVIDIA GPUs + CL_DEVICE_TYPE_GPU, "NVIDIA Corporation", { + { "GeForce GTX 480", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + { "Tesla K20m", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + { "Tesla K40m", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // AMD GPUs + CL_DEVICE_TYPE_GPU, "AMD", { + { "Tahiti", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // Intel GPUs + CL_DEVICE_TYPE_GPU, "Intel", { + { "Iris", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // Default + CL_DEVICE_TYPE_ALL, kDefault, { + { kDefault, { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvDouble = { + "Xgemv", Precision::kDouble, { + { // NVIDIA GPUs + CL_DEVICE_TYPE_GPU, "NVIDIA Corporation", { + { "GeForce GTX 480", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + { "Tesla K20m", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + { "Tesla K40m", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // AMD GPUs + CL_DEVICE_TYPE_GPU, "AMD", { + { "Tahiti", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // Intel GPUs + CL_DEVICE_TYPE_GPU, "Intel", { + } + }, + { // Default + CL_DEVICE_TYPE_ALL, kDefault, { + { kDefault, { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + } +}; +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvComplexSingle = { + "Xgemv", Precision::kComplexSingle, { + { // NVIDIA GPUs + CL_DEVICE_TYPE_GPU, "NVIDIA Corporation", { + { "GeForce GTX 480", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + { "Tesla K20m", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + { "Tesla K40m", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // AMD GPUs + CL_DEVICE_TYPE_GPU, "AMD", { + { "Tahiti", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // Intel GPUs + CL_DEVICE_TYPE_GPU, "Intel", { + { "Iris", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // Default + CL_DEVICE_TYPE_ALL, kDefault, { + { kDefault, { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvComplexDouble = { + "Xgemv", Precision::kComplexDouble, { + { // NVIDIA GPUs + CL_DEVICE_TYPE_GPU, "NVIDIA Corporation", { + { "GeForce GTX 480", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + { "Tesla K20m", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + { "Tesla K40m", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // AMD GPUs + CL_DEVICE_TYPE_GPU, "AMD", { + { "Tahiti", { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + { // Intel GPUs + CL_DEVICE_TYPE_GPU, "Intel", { + } + }, + { // Default + CL_DEVICE_TYPE_ALL, kDefault, { + { kDefault, { {"WGS",64}, {"WPT",1}, {"VW",1} } }, + } + }, + } +}; + +// ================================================================================================= +} // namespace clblast diff --git a/include/internal/routines/xgemv.h b/include/internal/routines/xgemv.h new file mode 100644 index 00000000..a3109036 --- /dev/null +++ b/include/internal/routines/xgemv.h @@ -0,0 +1,46 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren <www.cedricnugteren.nl> +// +// This file implements the Xgemv routine. The precision is implemented using a template argument. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XGEMV_H_ +#define CLBLAST_ROUTINES_XGEMV_H_ + +#include "internal/routine.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xgemv: public Routine { + public: + Xgemv(CommandQueue &queue, Event &event); + + // Templated-precision implementation of the routine + StatusCode DoGemv(const Layout layout, const Transpose a_transpose, + const size_t m, const size_t n, + const T alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const Buffer &x_buffer, const size_t x_offset, const size_t x_inc, + const T beta, + const Buffer &y_buffer, const size_t y_offset, const size_t y_inc); + + private: + // Static variable to get the precision + const static Precision precision_; +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XGEMV_H_ +#endif diff --git a/src/clblast.cc b/src/clblast.cc index 72de3b24..1d4d0621 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -20,6 +20,9 @@ // BLAS level-1 includes #include "internal/routines/xaxpy.h" +// BLAS level-2 includes +#include "internal/routines/xgemv.h" + // BLAS level-3 includes #include "internal/routines/xgemm.h" #include "internal/routines/xsymm.h" @@ -36,18 +39,18 @@ StatusCode Axpy(const size_t n, const T alpha, cl_command_queue* queue, cl_event* event) { auto queue_cpp = CommandQueue(*queue); auto event_cpp = Event(*event); - auto xaxpy = Xaxpy<T>(queue_cpp, event_cpp); + auto routine = Xaxpy<T>(queue_cpp, event_cpp); // Loads the kernel source-code as an include (C++11 raw string literal) std::string kernel_source = #include "kernels/xaxpy.opencl" - auto status = xaxpy.SetUp(kernel_source); + auto status = routine.SetUp(kernel_source); if (status != StatusCode::kSuccess) { return status; } // Runs the routine - return xaxpy.DoAxpy(n, alpha, - Buffer(x_buffer), x_offset, x_inc, - Buffer(y_buffer), y_offset, y_inc); + return routine.DoAxpy(n, alpha, + Buffer(x_buffer), x_offset, x_inc, + Buffer(y_buffer), y_offset, y_inc); } template StatusCode Axpy<float>(const size_t, const float, const cl_mem, const size_t, const size_t, @@ -69,22 +72,70 @@ template StatusCode Axpy<double2>(const size_t, const double2, // ================================================================================================= // BLAS level-2 (matrix-vector) routines +// GEMV +template <typename T> +StatusCode Gemv(const Layout layout, const Transpose transpose_a, + const size_t m, const size_t n, const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, const T beta, + cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + + auto queue_cpp = CommandQueue(*queue); + auto event_cpp = Event(*event); + auto routine = Xgemv<T>(queue_cpp, event_cpp); + + // Loads the kernel source-code as an include (C++11 raw string literal) + std::string kernel_source = + #include "kernels/xgemv.opencl" + auto status = routine.SetUp(kernel_source); + if (status != StatusCode::kSuccess) { return status; } + + // Runs the routine + return routine.DoGemv(layout, transpose_a, m, n, alpha, + Buffer(a_buffer), a_offset, a_ld, + Buffer(x_buffer), x_offset, x_inc, beta, + Buffer(y_buffer), y_offset, y_inc); +} +template StatusCode Gemv<float>(const Layout, const Transpose, + const size_t, const size_t, const float, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, const float, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); +template StatusCode Gemv<double>(const Layout, const Transpose, + const size_t, const size_t, const double, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, const double, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); +template StatusCode Gemv<float2>(const Layout, const Transpose, + const size_t, const size_t, const float2, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, const float2, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); +template StatusCode Gemv<double2>(const Layout, const Transpose, + const size_t, const size_t, const double2, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, const double2, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); + // ================================================================================================= // BLAS level-3 (matrix-matrix) routines // GEMM template <typename T> StatusCode Gemm(const Layout layout, const Transpose transpose_a, const Transpose transpose_b, - const size_t m, const size_t n, const size_t k, - const T alpha, + const size_t m, const size_t n, const size_t k, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, - const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, - const T beta, + const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, const T beta, cl_mem c_buffer, const size_t c_offset, const size_t c_ld, cl_command_queue* queue, cl_event* event) { auto queue_cpp = CommandQueue(*queue); auto event_cpp = Event(*event); - auto xgemm = Xgemm<T>(queue_cpp, event_cpp); + auto routine = Xgemm<T>(queue_cpp, event_cpp); // Loads the kernel source-code as an include (C++11 raw string literal) std::string common_source1 = @@ -97,50 +148,39 @@ StatusCode Gemm(const Layout layout, const Transpose transpose_a, const Transpos #include "kernels/padtranspose.opencl" std::string kernel_source = #include "kernels/xgemm.opencl" - auto status = xgemm.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + - kernel_source); + auto status = routine.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + + kernel_source); if (status != StatusCode::kSuccess) { return status; } // Runs the routine - return xgemm.DoGemm(layout, transpose_a, transpose_b, - m, n, k, - alpha, - Buffer(a_buffer), a_offset, a_ld, - Buffer(b_buffer), b_offset, b_ld, - beta, - Buffer(c_buffer), c_offset, c_ld); + return routine.DoGemm(layout, transpose_a, transpose_b, m, n, k, alpha, + Buffer(a_buffer), a_offset, a_ld, + Buffer(b_buffer), b_offset, b_ld, beta, + Buffer(c_buffer), c_offset, c_ld); } template StatusCode Gemm<float>(const Layout, const Transpose, const Transpose, - const size_t, const size_t, const size_t, - const float, - const cl_mem, const size_t, const size_t, + const size_t, const size_t, const size_t, const float, const cl_mem, const size_t, const size_t, - const float, + const cl_mem, const size_t, const size_t, const float, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); template StatusCode Gemm<double>(const Layout, const Transpose, const Transpose, - const size_t, const size_t, const size_t, - const double, - const cl_mem, const size_t, const size_t, + const size_t, const size_t, const size_t, const double, const cl_mem, const size_t, const size_t, - const double, + const cl_mem, const size_t, const size_t, const double, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); /* template StatusCode Gemm<float2>(const Layout, const Transpose, const Transpose, - const size_t, const size_t, const size_t, - const float2, + const size_t, const size_t, const size_t, const float2, const cl_mem, const size_t, const size_t, - const cl_mem, const size_t, const size_t, - const float2, + const cl_mem, const size_t, const size_t, const float2, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); template StatusCode Gemm<double2>(const Layout, const Transpose, const Transpose, - const size_t, const size_t, const size_t, - const double2, - const cl_mem, const size_t, const size_t, + const size_t, const size_t, const size_t, const double2, const cl_mem, const size_t, const size_t, - const double2, + const cl_mem, const size_t, const size_t, const double2, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); */ @@ -150,16 +190,14 @@ template StatusCode Gemm<double2>(const Layout, const Transpose, const Transpose // SYMM template <typename T> StatusCode Symm(const Layout layout, const Side side, const Triangle triangle, - const size_t m, const size_t n, - const T alpha, + const size_t m, const size_t n, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, - const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, - const T beta, + const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, const T beta, cl_mem c_buffer, const size_t c_offset, const size_t c_ld, cl_command_queue* queue, cl_event* event) { auto queue_cpp = CommandQueue(*queue); auto event_cpp = Event(*event); - auto xsymm = Xsymm<T>(queue_cpp, event_cpp); + auto routine = Xsymm<T>(queue_cpp, event_cpp); // Loads the kernel source-code as an include (C++11 raw string literal) std::string common_source1 = @@ -172,50 +210,39 @@ StatusCode Symm(const Layout layout, const Side side, const Triangle triangle, #include "kernels/padtranspose.opencl" std::string kernel_source = #include "kernels/xgemm.opencl" - auto status = xsymm.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + + auto status = routine.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + kernel_source); if (status != StatusCode::kSuccess) { return status; } // Runs the routine - return xsymm.DoSymm(layout, side, triangle, - m, n, - alpha, - Buffer(a_buffer), a_offset, a_ld, - Buffer(b_buffer), b_offset, b_ld, - beta, - Buffer(c_buffer), c_offset, c_ld); + return routine.DoSymm(layout, side, triangle, m, n, alpha, + Buffer(a_buffer), a_offset, a_ld, + Buffer(b_buffer), b_offset, b_ld, beta, + Buffer(c_buffer), c_offset, c_ld); } template StatusCode Symm<float>(const Layout, const Side, const Triangle, - const size_t, const size_t, - const float, + const size_t, const size_t, const float, const cl_mem, const size_t, const size_t, - const cl_mem, const size_t, const size_t, - const float, + const cl_mem, const size_t, const size_t, const float, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); template StatusCode Symm<double>(const Layout, const Side, const Triangle, - const size_t, const size_t, - const double, - const cl_mem, const size_t, const size_t, + const size_t, const size_t, const double, const cl_mem, const size_t, const size_t, - const double, + const cl_mem, const size_t, const size_t, const double, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); /* template StatusCode Symm<float2>(const Layout, const Side, const Triangle, - const size_t, const size_t, - const float2, + const size_t, const size_t, const float2, const cl_mem, const size_t, const size_t, - const cl_mem, const size_t, const size_t, - const float2, + const cl_mem, const size_t, const size_t, const float2, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); template StatusCode Symm<double2>(const Layout, const Side, const Triangle, - const size_t, const size_t, - const double2, - const cl_mem, const size_t, const size_t, + const size_t, const size_t, const double2, const cl_mem, const size_t, const size_t, - const double2, + const cl_mem, const size_t, const size_t, const double2, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); */ diff --git a/src/database.cc b/src/database.cc index beaa122b..4d9d844e 100644 --- a/src/database.cc +++ b/src/database.cc @@ -13,6 +13,7 @@ #include "internal/database.h" #include "internal/database/xaxpy.h" +#include "internal/database/xgemv.h" #include "internal/database/xgemm.h" #include "internal/database/copy.h" #include "internal/database/pad.h" @@ -27,6 +28,7 @@ namespace clblast { // Initializes the database const std::vector<Database::DatabaseEntry> Database::database = { XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble, + XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble, XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble, CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble, PadSingle, PadDouble, PadComplexSingle, PadComplexDouble, diff --git a/src/kernels/xgemv.opencl b/src/kernels/xgemv.opencl index c90bc26e..1d2ab435 100644 --- a/src/kernels/xgemv.opencl +++ b/src/kernels/xgemv.opencl @@ -34,7 +34,8 @@ R"( // The gemv kernel __attribute__((reqd_work_group_size(WGS, 1, 1))) __kernel void Xgemv(const int m, const int n, const real alpha, const real beta, - const __global real* restrict agm, + const int a_transposed, + const __global real* restrict agm, const int a_offset, const int a_ld, const __global real* restrict xgm, const int x_offset, const int x_inc, __global real* ygm, const int y_offset, const int y_inc) { @@ -45,8 +46,15 @@ __kernel void Xgemv(const int m, const int n, const real alpha, const real beta, // Loop over the elements of the matrix A real acc; SetToZero(acc); - for (int k=0; k<n; ++k) { - MultiplyAdd(acc, agm[id + m*k], xgm[k*x_inc + x_offset]); + if (a_transposed == 0) { + for (int k=0; k<n; ++k) { + MultiplyAdd(acc, agm[id + a_ld*k + a_offset], xgm[k*x_inc + x_offset]); + } + } + else { + for (int k=0; k<n; ++k) { + MultiplyAdd(acc, agm[k + a_ld*id + a_offset], xgm[k*x_inc + x_offset]); + } } AXPBY(ygm[id*y_inc + y_offset], alpha, acc, beta, ygm[id*y_inc + y_offset]); } diff --git a/src/routines/xgemv.cc b/src/routines/xgemv.cc new file mode 100644 index 00000000..67f9536e --- /dev/null +++ b/src/routines/xgemv.cc @@ -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 the Xgemv class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xgemv.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xgemv<float>::precision_ = Precision::kSingle; +template <> const Precision Xgemv<double>::precision_ = Precision::kDouble; +template <> const Precision Xgemv<float2>::precision_ = Precision::kComplexSingle; +template <> const Precision Xgemv<double2>::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xgemv<T>::Xgemv(CommandQueue &queue, Event &event): + Routine(queue, event, {"Xgemv"}, precision_) { +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xgemv<T>::DoGemv(const Layout layout, const Transpose a_transpose, + const size_t m, const size_t n, + const T alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const Buffer &x_buffer, const size_t x_offset, const size_t x_inc, + const T beta, + const Buffer &y_buffer, const size_t y_offset, const size_t y_inc) { + + // 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). + auto a_altlayout = (layout == Layout::kRowMajor); + auto a_one = (a_altlayout) ? n : m; + auto a_two = (a_altlayout) ? m : n; + + // Swap m and n if the matrix is transposed + auto a_transposed = (a_transpose == Transpose::kYes); + auto m_real = (a_transposed) ? n : m; + auto n_real = (a_transposed) ? m : n; + + // Determines whether the kernel needs to perform rotated access ('^' is the XOR operator) + auto a_rotated = a_transposed ^ a_altlayout; + + // 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(n_real, x_buffer, x_offset, x_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestVectorY(m_real, 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, "Xgemv"); + + // Sets the kernel arguments + kernel.SetArgument(0, static_cast<int>(m_real)); + kernel.SetArgument(1, static_cast<int>(n_real)); + kernel.SetArgument(2, alpha); + kernel.SetArgument(3, beta); + kernel.SetArgument(4, static_cast<int>(a_rotated)); + kernel.SetArgument(5, a_buffer()); + kernel.SetArgument(6, static_cast<int>(a_offset)); + kernel.SetArgument(7, static_cast<int>(a_ld)); + kernel.SetArgument(8, x_buffer()); + kernel.SetArgument(9, static_cast<int>(x_offset)); + kernel.SetArgument(10, static_cast<int>(x_inc)); + kernel.SetArgument(11, y_buffer()); + kernel.SetArgument(12, static_cast<int>(y_offset)); + kernel.SetArgument(13, static_cast<int>(y_inc)); + + // Launches the kernel + auto m_ceiled = Ceil(m_real, db_["WGS"]); + auto global = std::vector<size_t>{CeilDiv(m_ceiled, db_["WPT"])}; + auto local = std::vector<size_t>{db_["WGS"]}; + 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 Xgemv<float>; +template class Xgemv<double>; +template class Xgemv<float2>; +template class Xgemv<double2>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/tuning/xgemv.cc b/src/tuning/xgemv.cc index 1ee7c7bf..74bb77a5 100644 --- a/src/tuning/xgemv.cc +++ b/src/tuning/xgemv.cc @@ -55,7 +55,10 @@ void XgemvTune(const Arguments<T> &args, tuner.AddArgumentScalar(static_cast<int>(args.n)); tuner.AddArgumentScalar(args.alpha); tuner.AddArgumentScalar(args.beta); + tuner.AddArgumentScalar(0); tuner.AddArgumentInput(a_mat); + tuner.AddArgumentScalar(0); + tuner.AddArgumentScalar(static_cast<int>(args.n)); tuner.AddArgumentInput(x_vec); tuner.AddArgumentScalar(0); tuner.AddArgumentScalar(1); diff --git a/test/correctness/routines/xgemv.cc b/test/correctness/routines/xgemv.cc new file mode 100644 index 00000000..94ae147f --- /dev/null +++ b/test/correctness/routines/xgemv.cc @@ -0,0 +1,94 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under the MIT license. 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 tests for the Xgemv routine. It is based on the TestAXY class. +// +// ================================================================================================= + +#include "wrapper_clblas.h" +#include "correctness/testaxy.h" + +namespace clblast { +// ================================================================================================= + +// The correctness tester, containing the function calls to CLBlast and to clBLAS for comparison. +template <typename T> +void XgemvTest(int argc, char *argv[], const bool silent, const std::string &name) { + + // Creates the CLBlast lambda + auto clblast_lambda = [](const Arguments<T> &args, + const Buffer &a_mat, const Buffer &x_vec, const Buffer &y_vec, + CommandQueue &queue) -> StatusCode { + auto queue_plain = queue(); + auto event = cl_event{}; + return Gemv(args.layout, args.a_transpose, args.m, args.n, args.alpha, + a_mat(), args.a_offset, args.a_ld, + x_vec(), args.x_offset, args.x_inc, args.beta, + y_vec(), args.y_offset, args.y_inc, + &queue_plain, &event); + }; + + // Creates the clBLAS lambda (for comparison) + auto clblas_lambda = [](const Arguments<T> &args, + const Buffer &a_mat, const Buffer &x_vec, const Buffer &y_vec, + CommandQueue &queue) -> StatusCode { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXgemv(static_cast<clblasOrder>(args.layout), + static_cast<clblasTranspose>(args.a_transpose), + args.m, args.n, args.alpha, + a_mat(), args.a_offset, args.a_ld, + x_vec(), args.x_offset, args.x_inc, args.beta, + y_vec(), args.y_offset, args.y_inc, + 1, &queue_plain, 0, nullptr, &event); + return static_cast<StatusCode>(status); + }; + + // Selects the platform and device on which to test (command-line options) + auto help = std::string{"Options given/available:\n"}; + const auto platform_id = GetArgument(argc, argv, help, kArgPlatform, size_t{0}); + const auto device_id = GetArgument(argc, argv, help, kArgDevice, size_t{0}); + if (!silent) { fprintf(stdout, "\n* %s\n", help.c_str()); } + + // Initializes the other arguments relevant for this routine + auto args = Arguments<T>{}; + const auto options = std::vector<std::string>{kArgM, kArgN, kArgLayout, kArgATransp, + kArgALeadDim, kArgXInc, kArgYInc, + kArgAOffset, kArgXOffset, kArgYOffset}; + + // Creates a tester + TestAXY<T> tester{platform_id, device_id, name, options, clblast_lambda, clblas_lambda}; + + // Loops over the test-cases from a data-layout point of view + for (auto &layout: {Layout::kRowMajor, Layout::kColMajor}) { + args.layout = layout; + for (auto &a_transpose: {Transpose::kNo, Transpose::kYes}) { + args.a_transpose = a_transpose; + const auto case_name = ToString(layout)+" "+ToString(a_transpose); + + // Runs the tests + tester.TestRegular(args, case_name); + tester.TestInvalidBufferSizes(args, case_name); + } + } +} + +// ================================================================================================= +} // namespace clblast + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::XgemvTest<float>(argc, argv, false, "SGEMV"); + //clblast::XgemvTest<double>(argc, argv, true, "DGEMV"); + //clblast::XgemvTest<clblast::float2>(argc, argv, true, "CGEMV"); + //clblast::XgemvTest<clblast::double2>(argc, argv, true, "ZGEMV"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/testaxy.cc b/test/correctness/testaxy.cc index 1e01a0e8..ed0b06ab 100644 --- a/test/correctness/testaxy.cc +++ b/test/correctness/testaxy.cc @@ -49,12 +49,6 @@ template <typename T> void TestAXY<T>::TestRegular(Arguments<T> &args, const std::string &name) { TestStart("regular behaviour", name); - // Computes whether or not the matrix is transposed. Note that we assume a default of - // column-major and no-transpose. If one of them is different (but not both), then rotated - // is considered true. - auto a_rotated = (args.layout == Layout::kColMajor && args.a_transpose != Transpose::kNo) || - (args.layout == Layout::kRowMajor && args.a_transpose == Transpose::kNo); - // Iterates over the dimension for the matrix and vectors for (auto &m: kMatrixVectorDims) { args.m = m; @@ -62,7 +56,12 @@ void TestAXY<T>::TestRegular(Arguments<T> &args, const std::string &name) { args.n = n; // Computes the second dimension of the matrix taking the rotation into account - auto a_two = (a_rotated) ? m : n; + auto a_two = (args.layout == Layout::kRowMajor) ? args.m : args.n; + + // Computes the vector sizes in case the matrix is transposed + auto a_transposed = (args.a_transpose == Transpose::kYes); + auto m_real = (a_transposed) ? n : m; + auto n_real = (a_transposed) ? m : n; // Iterates over the leading-dimension values and the offsets of the matrix for (auto &a_ld: kMatrixVectorDims) { @@ -82,8 +81,8 @@ void TestAXY<T>::TestRegular(Arguments<T> &args, const std::string &name) { // Computes the buffer sizes auto a_size = a_two * a_ld + a_offset; - auto x_size = n * x_inc + x_offset; - auto y_size = n * y_inc + y_offset; + auto x_size = n_real * x_inc + x_offset; + auto y_size = m_real * y_inc + y_offset; if (a_size < 1 || x_size < 1 || y_size < 1) { continue; } // Creates the OpenCL buffers @@ -124,15 +123,15 @@ void TestAXY<T>::TestRegular(Arguments<T> &args, const std::string &name) { // Checks for differences in the output auto errors = size_t{0}; - for (auto idn=size_t{0}; idn<n; ++idn) { - auto index = idn*y_inc + y_offset; + for (auto idm=size_t{0}; idm<m_real; ++idm) { + auto index = idm*y_inc + y_offset; if (!TestSimilarity(r_result[index], s_result[index], kErrorMargin)) { errors++; } } // Tests the error count (should be zero) - TestErrorCount(errors, n, args); + TestErrorCount(errors, m_real, args); } } } @@ -158,6 +157,9 @@ void TestAXY<T>::TestInvalidBufferSizes(Arguments<T> &args, const std::string &n args.m = kBufferSize; args.n = kBufferSize; args.a_ld = kBufferSize; + args.a_offset = 0; + args.x_offset = 0; + args.y_offset = 0; // Iterates over test buffer sizes const std::vector<size_t> kMatrixSizes = {0, kBufferSize*kBufferSize-1, kBufferSize*kBufferSize}; diff --git a/test/performance/client.cc b/test/performance/client.cc index ddaea0e1..3b07970c 100644 --- a/test/performance/client.cc +++ b/test/performance/client.cc @@ -26,8 +26,12 @@ template <typename T> void ClientXY(int argc, char *argv[], Routine2<T> client_routine, const std::vector<std::string> &options) { + // Function to determine how to find the default value of the leading dimension of matrix A. + // Note: this is not relevant for this client but given anyway. + auto default_ld_a = [](const Arguments<T> args) { return args.n; }; + // Simple command line argument parser with defaults - auto args = ParseArguments<T>(argc, argv, options); + auto args = ParseArguments<T>(argc, argv, options, default_ld_a); if (args.print_help) { return; } // Prints the header of the output table @@ -81,13 +85,94 @@ template void ClientXY<double2>(int, char **, Routine2<double2>, const std::vect // ================================================================================================= +// This is the matrix-vector-vector variant of the set-up/tear-down client routine. +template <typename T> +void ClientAXY(int argc, char *argv[], Routine3<T> client_routine, + const std::vector<std::string> &options) { + + // Function to determine how to find the default value of the leading dimension of matrix A + auto default_ld_a = [](const Arguments<T> args) { return args.n; }; + + // Simple command line argument parser with defaults + auto args = ParseArguments<T>(argc, argv, options, default_ld_a); + if (args.print_help) { return; } + + // Prints the header of the output table + PrintTableHeader(args.silent, options); + + // Initializes OpenCL and the libraries + auto platform = Platform(args.platform_id); + auto device = Device(platform, kDeviceType, args.device_id); + auto context = Context(device); + auto queue = CommandQueue(context, device); + if (args.compare_clblas) { clblasSetup(); } + + // Iterates over all "num_step" values jumping by "step" each time + auto s = size_t{0}; + while(true) { + + // Computes the second dimension of the matrix taking the rotation into account + auto a_two = (args.layout == Layout::kRowMajor) ? args.m : args.n; + + // Computes the vector sizes in case the matrix is transposed + auto a_transposed = (args.a_transpose == Transpose::kYes); + auto m_real = (a_transposed) ? args.n : args.m; + auto n_real = (a_transposed) ? args.m : args.n; + + // Computes the data sizes + auto a_size = a_two * args.a_ld + args.a_offset; + auto x_size = n_real*args.x_inc + args.x_offset; + auto y_size = m_real*args.y_inc + args.y_offset; + + // Populates input host vectors with random data + std::vector<T> a_source(a_size); + std::vector<T> x_source(x_size); + std::vector<T> y_source(y_size); + PopulateVector(a_source); + PopulateVector(x_source); + PopulateVector(y_source); + + // Creates the vectors on the device + auto a_buffer = Buffer(context, CL_MEM_READ_WRITE, a_size*sizeof(T)); + auto x_buffer = Buffer(context, CL_MEM_READ_WRITE, x_size*sizeof(T)); + auto y_buffer = Buffer(context, CL_MEM_READ_WRITE, y_size*sizeof(T)); + a_buffer.WriteBuffer(queue, a_size*sizeof(T), a_source); + x_buffer.WriteBuffer(queue, x_size*sizeof(T), x_source); + y_buffer.WriteBuffer(queue, y_size*sizeof(T), y_source); + + // Runs the routine-specific code + client_routine(args, a_buffer, x_buffer, y_buffer, queue); + + // Makes the jump to the next step + ++s; + if (s >= args.num_steps) { break; } + args.m += args.step; + args.n += args.step; + args.a_ld += args.step; + } + + // Cleans-up and returns + if (args.compare_clblas) { clblasTeardown(); } +} + +// Compiles the above function +template void ClientAXY<float>(int, char **, Routine3<float>, const std::vector<std::string>&); +template void ClientAXY<double>(int, char **, Routine3<double>, const std::vector<std::string>&); +template void ClientAXY<float2>(int, char **, Routine3<float2>, const std::vector<std::string>&); +template void ClientAXY<double2>(int, char **, Routine3<double2>, const std::vector<std::string>&); + +// ================================================================================================= + // This is the matrix-matrix-matrix variant of the set-up/tear-down client routine. template <typename T> void ClientABC(int argc, char *argv[], Routine3<T> client_routine, const std::vector<std::string> &options) { + // Function to determine how to find the default value of the leading dimension of matrix A + auto default_ld_a = [](const Arguments<T> args) { return args.m; }; + // Simple command line argument parser with defaults - auto args = ParseArguments<T>(argc, argv, options); + auto args = ParseArguments<T>(argc, argv, options, default_ld_a); if (args.print_help) { return; } // Prints the header of the output table @@ -167,7 +252,8 @@ template void ClientABC<double2>(int, char **, Routine3<double2>, const std::vec // applicable, but are searched for anyway to be able to create one common argument parser. All // arguments have a default value in case they are not found. template <typename T> -Arguments<T> ParseArguments(int argc, char *argv[], const std::vector<std::string> &options) { +Arguments<T> ParseArguments(int argc, char *argv[], const std::vector<std::string> &options, + const std::function<size_t(const Arguments<T>)> default_ld_a) { auto args = Arguments<T>{}; auto help = std::string{"Options given/available:\n"}; @@ -193,7 +279,7 @@ Arguments<T> ParseArguments(int argc, char *argv[], const std::vector<std::strin if (o == kArgYOffset) { args.y_offset = GetArgument(argc, argv, help, kArgYOffset, size_t{0}); } // Matrix arguments - if (o == kArgALeadDim) { args.a_ld = GetArgument(argc, argv, help, kArgALeadDim, args.k); } + if (o == kArgALeadDim) { args.a_ld = GetArgument(argc, argv, help, kArgALeadDim, default_ld_a(args)); } if (o == kArgBLeadDim) { args.b_ld = GetArgument(argc, argv, help, kArgBLeadDim, args.n); } if (o == kArgCLeadDim) { args.c_ld = GetArgument(argc, argv, help, kArgCLeadDim, args.n); } if (o == kArgAOffset) { args.a_offset = GetArgument(argc, argv, help, kArgAOffset, size_t{0}); } diff --git a/test/performance/client.h b/test/performance/client.h index 2b9991fe..5125844a 100644 --- a/test/performance/client.h +++ b/test/performance/client.h @@ -49,6 +49,9 @@ template <typename T> void ClientXY(int argc, char *argv[], Routine2<T> client_routine, const std::vector<std::string> &options); template <typename T> +void ClientAXY(int argc, char *argv[], Routine3<T> client_routine, + const std::vector<std::string> &options); +template <typename T> void ClientABC(int argc, char *argv[], Routine3<T> client_routine, const std::vector<std::string> &options); @@ -57,7 +60,8 @@ void ClientABC(int argc, char *argv[], Routine3<T> client_routine, // Parses all command-line arguments, filling in the arguments structure. If no command-line // argument is given for a particular argument, it is filled in with a default value. template <typename T> -Arguments<T> ParseArguments(int argc, char *argv[], const std::vector<std::string> &options); +Arguments<T> ParseArguments(int argc, char *argv[], const std::vector<std::string> &options, + const std::function<size_t(const Arguments<T>)> default_ld_a); // Retrieves only the precision command-line argument, since the above function is templated based // on the precision diff --git a/test/performance/routines/xgemv.cc b/test/performance/routines/xgemv.cc new file mode 100644 index 00000000..43222396 --- /dev/null +++ b/test/performance/routines/xgemv.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 Xgemv command-line interface tester. +// +// ================================================================================================= + +#include <string> +#include <vector> +#include <exception> + +#include "wrapper_clblas.h" +#include "performance/client.h" + +namespace clblast { +// ================================================================================================= + +// The client, used for performance testing. It contains the function calls to CLBlast and to other +// libraries to compare against. +template <typename T> +void PerformanceXgemv(const Arguments<T> &args, + const Buffer &a_mat, const Buffer &x_vec, const Buffer &y_vec, + CommandQueue &queue) { + + // Creates the CLBlast lambda + auto clblast_lambda = [&args, &a_mat, &x_vec, &y_vec, &queue]() { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = Gemv(args.layout, args.a_transpose, args.m, args.n, args.alpha, + a_mat(), args.a_offset, args.a_ld, + x_vec(), args.x_offset, args.x_inc, args.beta, + y_vec(), args.y_offset, args.y_inc, + &queue_plain, &event); + clWaitForEvents(1, &event); + if (status != StatusCode::kSuccess) { + throw std::runtime_error("CLBlast error: "+ToString(static_cast<int>(status))); + } + }; + + // Creates the clBLAS lambda (for comparison) + auto clblas_lambda = [&args, &a_mat, &x_vec, &y_vec, &queue]() { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXgemv(static_cast<clblasOrder>(args.layout), + static_cast<clblasTranspose>(args.a_transpose), + args.m, args.n, args.alpha, + a_mat(), args.a_offset, args.a_ld, + x_vec(), args.x_offset, args.x_inc, args.beta, + y_vec(), args.y_offset, args.y_inc, + 1, &queue_plain, 0, nullptr, &event); + clWaitForEvents(1, &event); + if (status != CL_SUCCESS) { + throw std::runtime_error("clBLAS error: "+ToString(static_cast<int>(status))); + } + }; + + // Runs the routines and collect the timings + auto ms_clblast = TimedExecution(args.num_runs, clblast_lambda); + auto ms_clblas = TimedExecution(args.num_runs, clblas_lambda); + + // Prints the performance of both libraries + const auto flops = 2 * args.m * args.n; + const auto bytes = (args.m*args.n + 2*args.m + args.n) * sizeof(T); + const auto output_ints = std::vector<size_t>{args.m, args.n, + static_cast<size_t>(args.layout), + static_cast<size_t>(args.a_transpose), + args.a_ld, args.x_inc, args.y_inc, + args.a_offset, args.x_offset, args.y_offset}; + const auto output_strings = std::vector<std::string>{ToString(args.alpha), + ToString(args.beta)}; + PrintTableRow(output_ints, output_strings, args.no_abbrv, + ms_clblast, ms_clblas, flops, bytes); +} + +// ================================================================================================= + +// Main function which calls the common client code with the routine-specific function as argument. +void ClientXgemv(int argc, char *argv[]) { + const auto o = std::vector<std::string>{kArgM, kArgN, kArgLayout, kArgATransp, + kArgALeadDim, kArgXInc, kArgYInc, + kArgAOffset, kArgXOffset, kArgYOffset, + kArgAlpha, kArgBeta}; + switch(GetPrecision(argc, argv)) { + case Precision::kHalf: throw std::runtime_error("Unsupported precision mode"); + case Precision::kSingle: ClientAXY<float>(argc, argv, PerformanceXgemv<float>, o); break; + case Precision::kDouble: ClientAXY<double>(argc, argv, PerformanceXgemv<double>, o); break; + case Precision::kComplexSingle: ClientAXY<float2>(argc, argv, PerformanceXgemv<float2>, o); break; + case Precision::kComplexDouble: ClientAXY<double2>(argc, argv, PerformanceXgemv<double2>, o); break; + } +} + +// ================================================================================================= +} // namespace clblast + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::ClientXgemv(argc, argv); + return 0; +} + +// ================================================================================================= diff --git a/test/wrapper_clblas.h b/test/wrapper_clblas.h index 7c71fcaa..093a8742 100644 --- a/test/wrapper_clblas.h +++ b/test/wrapper_clblas.h @@ -74,6 +74,64 @@ clblasStatus clblasXaxpy( // ================================================================================================= // BLAS level-2 (matrix-vector) routines +// Calls {clblasSgemv, clblasDgemv, clblasCgemv, clblasZgemv} with the arguments forwarded. +clblasStatus clblasXgemv( + clblasOrder layout, clblasTranspose tran_a, size_t m, size_t n, float alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem x_vec, size_t x_offset, size_t x_inc, float beta, + const cl_mem y_vec, size_t y_offset, size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasSgemv(layout, tran_a, m, n, alpha, + a_mat, a_offset, a_ld, + x_vec, x_offset, static_cast<int>(x_inc), beta, + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXgemv( + clblasOrder layout, clblasTranspose tran_a, size_t m, size_t n, double alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem x_vec, size_t x_offset, size_t x_inc, double beta, + const cl_mem y_vec, size_t y_offset, size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasDgemv(layout, tran_a, m, n, alpha, + a_mat, a_offset, a_ld, + x_vec, x_offset, static_cast<int>(x_inc), beta, + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXgemv( + clblasOrder layout, clblasTranspose tran_a, size_t m, size_t n, float2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem x_vec, size_t x_offset, size_t x_inc, float2 beta, + const cl_mem y_vec, size_t y_offset, size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto cl_alpha = cl_float2{{alpha.real(), alpha.imag()}}; + auto cl_beta = cl_float2{{beta.real(), beta.imag()}}; + return clblasCgemv(layout, tran_a, m, n, cl_alpha, + a_mat, a_offset, a_ld, + x_vec, x_offset, static_cast<int>(x_inc), cl_beta, + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXgemv( + clblasOrder layout, clblasTranspose tran_a, size_t m, size_t n, double2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem x_vec, size_t x_offset, size_t x_inc, double2 beta, + const cl_mem y_vec, size_t y_offset, size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto cl_alpha = cl_double2{{alpha.real(), alpha.imag()}}; + auto cl_beta = cl_double2{{beta.real(), beta.imag()}}; + return clblasZgemv(layout, tran_a, m, n, cl_alpha, + a_mat, a_offset, a_ld, + x_vec, x_offset, static_cast<int>(x_inc), cl_beta, + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} + // ================================================================================================= // BLAS level-3 (matrix-matrix) routines |