diff options
| author | CNugteren <web@cedricnugteren.nl> | 2015-08-22 17:11:20 +0200 |
|---|---|---|
| committer | CNugteren <web@cedricnugteren.nl> | 2015-08-22 17:11:20 +0200 |
| commit | ff0c54c3865b45eff807315262e73d3f01cb19c3 (patch) | |
| tree | 839e9def73fb068f988b07e1e879ecce48d884c8 | |
| parent | 75517353d505de1d3979866060261a666aebfd36 (diff) | |
Added the XSWAP, XSCAL and XCOPY level-1 routines
29 files changed, 1811 insertions, 58 deletions
@@ -1,6 +1,9 @@ Development version (next release) -- +- Added level-1 routines: + * SSWAP/DSWAP/CSWAP/ZSWAP + * SSCAL/DSCAL/CSCAL/ZSCAL + * SCOPY/DCOPY/CCOPY/ZCOPY Version 0.4.0 - Now using the Claduc C++11 interface to OpenCL diff --git a/CMakeLists.txt b/CMakeLists.txt index c9a398a7..ba89bf91 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -105,7 +105,7 @@ include_directories(${clblast_SOURCE_DIR}/include ${OPENCL_INCLUDE_DIRS}) set(KERNELS copy pad transpose padtranspose xaxpy xgemv xgemm) set(SAMPLE_PROGRAMS_CPP sgemm) set(SAMPLE_PROGRAMS_C sgemm) -set(LEVEL1_ROUTINES xaxpy) +set(LEVEL1_ROUTINES xswap xscal xcopy xaxpy) set(LEVEL2_ROUTINES xgemv xhemv xsymv) set(LEVEL3_ROUTINES xgemm xsymm xhemm xsyrk xherk xsyr2k xher2k xtrmm) set(ROUTINES ${LEVEL1_ROUTINES} ${LEVEL2_ROUTINES} ${LEVEL3_ROUTINES}) @@ -138,9 +138,9 @@ CLBlast is in active development and currently does not support the full set of | xROTMG | | | - | - | | | xROT | | | - | - | | | xROTM | | | - | - | | -| xSWAP | | | | | | -| xSCAL | | | | | +CS +ZD | -| xCOPY | | | | | | +| xSWAP | ✔ | ✔ | ✔ | ✔ | | +| xSCAL | ✔ | ✔ | ✔ | ✔ | +CS +ZD | +| xCOPY | ✔ | ✔ | ✔ | ✔ | | | xAXPY | ✔ | ✔ | ✔ | ✔ | | | xDOT | | | - | - | +DS | | xDOTU | - | - | | | | diff --git a/include/clblast.h b/include/clblast.h index bd0f161c..326c9ec2 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -83,16 +83,40 @@ enum class Precision { kHalf = 16, kSingle = 32, kDouble = 64, // ================================================================================================= // BLAS level-1 (vector-vector) routines +// ================================================================================================= + +// Swap two vectors: SSWAP/DSWAP/CSWAP/ZSWAP +template <typename T> +StatusCode Swap(const size_t n, + 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); + +// Vector scaling: SSCAL/DSCAL/CSCAL/ZSCAL +template <typename T> +StatusCode Scal(const size_t n, + const T alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event); -// Templated-precision vector-times-constant plus vector: SAXPY/DAXPY/CAXPY/ZAXPY +// Vector copy: SCOPY/DCOPY/CCOPY/ZCOPY template <typename T> -StatusCode Axpy(const size_t n, const T alpha, +StatusCode Copy(const size_t n, + 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); + +// Vector-times-constant plus vector: SAXPY/DAXPY/CAXPY/ZAXPY +template <typename T> +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); // ================================================================================================= // BLAS level-2 (matrix-vector) routines +// ================================================================================================= // Templated-precision generalized matrix-vector multiplication: SGEMV/DGEMV/CGEMV/ZGEMV template <typename T> @@ -129,6 +153,7 @@ StatusCode Symv(const Layout layout, const Triangle triangle, // ================================================================================================= // BLAS level-3 (matrix-matrix) routines +// ================================================================================================= // Templated-precision generalized matrix-matrix multiplication: SGEMM/DGEMM/CGEMM/ZGEMM template <typename T> diff --git a/include/clblast_c.h b/include/clblast_c.h index c25e5880..766570e6 100644 --- a/include/clblast_c.h +++ b/include/clblast_c.h @@ -81,6 +81,60 @@ typedef enum Precision_ { kHalf = 16, kSingle = 32, kDouble = 64, // BLAS level-1 (vector-vector) routines // ================================================================================================= +// Swap two vectors: SSWAP/DSWAP/CSWAP/ZSWAP +StatusCode CLBlastSswap(const size_t n, + 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); +StatusCode CLBlastDswap(const size_t n, + 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); +StatusCode CLBlastCswap(const size_t n, + 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); +StatusCode CLBlastZswap(const size_t n, + 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); + +// Vector scaling: SSCAL/DSCAL/CSCAL/ZSCAL +StatusCode CLBlastSscal(const size_t n, + const float alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event); +StatusCode CLBlastDscal(const size_t n, + const double alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event); +StatusCode CLBlastCscal(const size_t n, + const cl_float2 alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event); +StatusCode CLBlastZscal(const size_t n, + const cl_double2 alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event); + +// Vector copy: SCOPY/DCOPY/CCOPY/ZCOPY +StatusCode CLBlastScopy(const size_t n, + 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); +StatusCode CLBlastDcopy(const size_t n, + 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); +StatusCode CLBlastCcopy(const size_t n, + 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); +StatusCode CLBlastZcopy(const size_t n, + 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); + // Vector-times-constant plus vector: SAXPY/DAXPY/CAXPY/ZAXPY StatusCode CLBlastSaxpy(const size_t n, const float alpha, diff --git a/include/internal/routines/level1/xcopy.h b/include/internal/routines/level1/xcopy.h new file mode 100644 index 00000000..c71583c5 --- /dev/null +++ b/include/internal/routines/level1/xcopy.h @@ -0,0 +1,54 @@ + +// ================================================================================================= +// 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 Xcopy routine. The precision is implemented using a template argument. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XCOPY_H_ +#define CLBLAST_ROUTINES_XCOPY_H_ + +#include "internal/routine.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xcopy: public Routine<T> { + public: + + // Members and methods from the base class + using Routine<T>::db_; + using Routine<T>::source_string_; + using Routine<T>::queue_; + using Routine<T>::GetProgramFromCache; + using Routine<T>::TestVectorX; + using Routine<T>::TestVectorY; + using Routine<T>::RunKernel; + using Routine<T>::ErrorIn; + + // Constructor + Xcopy(Queue &queue, Event &event); + + // Templated-precision implementation of the routine + StatusCode DoCopy(const size_t n, + 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); + + private: + // Static variable to get the precision + const static Precision precision_; +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XCOPY_H_ +#endif diff --git a/include/internal/routines/level1/xscal.h b/include/internal/routines/level1/xscal.h new file mode 100644 index 00000000..0aa6059d --- /dev/null +++ b/include/internal/routines/level1/xscal.h @@ -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 Xscal routine. The precision is implemented using a template argument. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XSCAL_H_ +#define CLBLAST_ROUTINES_XSCAL_H_ + +#include "internal/routine.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xscal: public Routine<T> { + public: + + // Members and methods from the base class + using Routine<T>::db_; + using Routine<T>::source_string_; + using Routine<T>::queue_; + using Routine<T>::GetProgramFromCache; + using Routine<T>::TestVectorX; + using Routine<T>::RunKernel; + using Routine<T>::ErrorIn; + + // Constructor + Xscal(Queue &queue, Event &event); + + // Templated-precision implementation of the routine + StatusCode DoScal(const size_t n, const T alpha, + const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc); + + private: + // Static variable to get the precision + const static Precision precision_; +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XSCAL_H_ +#endif diff --git a/include/internal/routines/level1/xswap.h b/include/internal/routines/level1/xswap.h new file mode 100644 index 00000000..3dabc62c --- /dev/null +++ b/include/internal/routines/level1/xswap.h @@ -0,0 +1,54 @@ + +// ================================================================================================= +// 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 Xswap routine. The precision is implemented using a template argument. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XSWAP_H_ +#define CLBLAST_ROUTINES_XSWAP_H_ + +#include "internal/routine.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xswap: public Routine<T> { + public: + + // Members and methods from the base class + using Routine<T>::db_; + using Routine<T>::source_string_; + using Routine<T>::queue_; + using Routine<T>::GetProgramFromCache; + using Routine<T>::TestVectorX; + using Routine<T>::TestVectorY; + using Routine<T>::RunKernel; + using Routine<T>::ErrorIn; + + // Constructor + Xswap(Queue &queue, Event &event); + + // Templated-precision implementation of the routine + StatusCode DoSwap(const size_t n, + 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); + + private: + // Static variable to get the precision + const static Precision precision_; +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XSWAP_H_ +#endif diff --git a/src/clblast.cc b/src/clblast.cc index 12c7b880..c99ad7b1 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -18,6 +18,9 @@ #include "clblast.h" // BLAS level-1 includes +#include "internal/routines/level1/xswap.h" +#include "internal/routines/level1/xscal.h" +#include "internal/routines/level1/xcopy.h" #include "internal/routines/level1/xaxpy.h" // BLAS level-2 includes @@ -40,41 +43,139 @@ namespace clblast { // BLAS level-1 (vector-vector) routines // ================================================================================================= +// SWAP +template <typename T> +StatusCode Swap(const size_t n, + 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) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xswap<T>(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoSwap(n, + Buffer<T>(x_buffer), x_offset, x_inc, + Buffer<T>(y_buffer), y_offset, y_inc); +} +template StatusCode Swap<float>(const size_t, + cl_mem, const size_t, const size_t, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Swap<double>(const size_t, + cl_mem, const size_t, const size_t, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Swap<float2>(const size_t, + cl_mem, const size_t, const size_t, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Swap<double2>(const size_t, + cl_mem, const size_t, const size_t, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); + +// SCAL +template <typename T> +StatusCode Scal(const size_t n, + const T alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xscal<T>(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoScal(n, + alpha, + Buffer<T>(x_buffer), x_offset, x_inc); +} +template StatusCode Scal<float>(const size_t, + const float, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Scal<double>(const size_t, + const double, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Scal<float2>(const size_t, + const float2, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Scal<double2>(const size_t, + const double2, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); + +// COPY +template <typename T> +StatusCode Copy(const size_t n, + 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) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xcopy<T>(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoCopy(n, + Buffer<T>(x_buffer), x_offset, x_inc, + Buffer<T>(y_buffer), y_offset, y_inc); +} +template StatusCode Copy<float>(const size_t, + const cl_mem, const size_t, const size_t, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Copy<double>(const size_t, + const cl_mem, const size_t, const size_t, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Copy<float2>(const size_t, + const cl_mem, const size_t, const size_t, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Copy<double2>(const size_t, + const cl_mem, const size_t, const size_t, + cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); + // AXPY template <typename T> -StatusCode Axpy(const size_t n, 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) { auto queue_cpp = Queue(*queue); auto event_cpp = Event(*event); auto routine = Xaxpy<T>(queue_cpp, event_cpp); - - // Compiles the routine's device kernels auto status = routine.SetUp(); if (status != StatusCode::kSuccess) { return status; } - - // Runs the routine - return routine.DoAxpy(n, alpha, + return routine.DoAxpy(n, + alpha, Buffer<T>(x_buffer), x_offset, x_inc, Buffer<T>(y_buffer), y_offset, y_inc); } -template StatusCode Axpy<float>(const size_t, const float, +template StatusCode Axpy<float>(const size_t, + const float, const cl_mem, const size_t, const size_t, cl_mem, const size_t, const size_t, - cl_command_queue*, cl_event*); -template StatusCode Axpy<double>(const size_t, const double, + cl_command_queue* queue, cl_event* event); +template StatusCode Axpy<double>(const size_t, + const double, const cl_mem, const size_t, const size_t, cl_mem, const size_t, const size_t, - cl_command_queue*, cl_event*); -template StatusCode Axpy<float2>(const size_t, const float2, + cl_command_queue* queue, cl_event* event); +template StatusCode Axpy<float2>(const size_t, + const float2, const cl_mem, const size_t, const size_t, cl_mem, const size_t, const size_t, - cl_command_queue*, cl_event*); -template StatusCode Axpy<double2>(const size_t, const double2, + cl_command_queue* queue, cl_event* event); +template StatusCode Axpy<double2>(const size_t, + const double2, const cl_mem, const size_t, const size_t, cl_mem, const size_t, const size_t, - cl_command_queue*, cl_event*); + cl_command_queue* queue, cl_event* event); // ================================================================================================= // BLAS level-2 (matrix-vector) routines diff --git a/src/clblast_c.cc b/src/clblast_c.cc index 3b437aff..ab3994fb 100644 --- a/src/clblast_c.cc +++ b/src/clblast_c.cc @@ -19,10 +19,140 @@ extern "C" { #include "clblast.h" #include "internal/utilities.h" +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + // ================================================================================================= // BLAS level-1 (vector-vector) routines // ================================================================================================= +// SWAP +StatusCode CLBlastSswap(const size_t n, + 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) { + auto status = clblast::Swap<float>(n, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastDswap(const size_t n, + 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) { + auto status = clblast::Swap<double>(n, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastCswap(const size_t n, + 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) { + auto status = clblast::Swap<float2>(n, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastZswap(const size_t n, + 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) { + auto status = clblast::Swap<double2>(n, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} + +// SCAL +StatusCode CLBlastSscal(const size_t n, + const float alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Scal(n, + alpha, + x_buffer, x_offset, x_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastDscal(const size_t n, + const double alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Scal(n, + alpha, + x_buffer, x_offset, x_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastCscal(const size_t n, + const cl_float2 alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Scal(n, + float2{alpha.s[0], alpha.s[1]}, + x_buffer, x_offset, x_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastZscal(const size_t n, + const cl_double2 alpha, + cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Scal(n, + double2{alpha.s[0], alpha.s[1]}, + x_buffer, x_offset, x_inc, + queue, event); + return static_cast<StatusCode>(status); +} + +// COPY +StatusCode CLBlastScopy(const size_t n, + 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) { + auto status = clblast::Copy<float>(n, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastDcopy(const size_t n, + 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) { + auto status = clblast::Copy<double>(n, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastCcopy(const size_t n, + 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) { + auto status = clblast::Copy<float2>(n, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastZcopy(const size_t n, + 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) { + auto status = clblast::Copy<double2>(n, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} + // AXPY StatusCode CLBlastSaxpy(const size_t n, const float alpha, @@ -54,7 +184,7 @@ StatusCode CLBlastCaxpy(const size_t n, cl_mem y_buffer, const size_t y_offset, const size_t y_inc, cl_command_queue* queue, cl_event* event) { auto status = clblast::Axpy(n, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, x_buffer, x_offset, x_inc, y_buffer, y_offset, y_inc, queue, event); @@ -66,7 +196,7 @@ StatusCode CLBlastZaxpy(const size_t n, cl_mem y_buffer, const size_t y_offset, const size_t y_inc, cl_command_queue* queue, cl_event* event) { auto status = clblast::Axpy(n, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, x_buffer, x_offset, x_inc, y_buffer, y_offset, y_inc, queue, event); @@ -127,10 +257,10 @@ StatusCode CLBlastCgemv(const Layout layout, const Transpose a_transpose, auto status = clblast::Gemv(static_cast<clblast::Layout>(layout), static_cast<clblast::Transpose>(a_transpose), m, n, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, x_buffer, x_offset, x_inc, - clblast::float2{beta.s[0], beta.s[1]}, + float2{beta.s[0], beta.s[1]}, y_buffer, y_offset, y_inc, queue, event); return static_cast<StatusCode>(status); @@ -146,10 +276,10 @@ StatusCode CLBlastZgemv(const Layout layout, const Transpose a_transpose, auto status = clblast::Gemv(static_cast<clblast::Layout>(layout), static_cast<clblast::Transpose>(a_transpose), m, n, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, x_buffer, x_offset, x_inc, - clblast::double2{beta.s[0], beta.s[1]}, + double2{beta.s[0], beta.s[1]}, y_buffer, y_offset, y_inc, queue, event); return static_cast<StatusCode>(status); @@ -167,10 +297,10 @@ StatusCode CLBlastChemv(const Layout layout, const Triangle triangle, auto status = clblast::Hemv(static_cast<clblast::Layout>(layout), static_cast<clblast::Triangle>(triangle), n, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, x_buffer, x_offset, x_inc, - clblast::float2{beta.s[0], beta.s[1]}, + float2{beta.s[0], beta.s[1]}, y_buffer, y_offset, y_inc, queue, event); return static_cast<StatusCode>(status); @@ -186,10 +316,10 @@ StatusCode CLBlastZhemv(const Layout layout, const Triangle triangle, auto status = clblast::Hemv(static_cast<clblast::Layout>(layout), static_cast<clblast::Triangle>(triangle), n, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, x_buffer, x_offset, x_inc, - clblast::double2{beta.s[0], beta.s[1]}, + double2{beta.s[0], beta.s[1]}, y_buffer, y_offset, y_inc, queue, event); return static_cast<StatusCode>(status); @@ -292,10 +422,10 @@ StatusCode CLBlastCgemm(const Layout layout, const Transpose a_transpose, const static_cast<clblast::Transpose>(a_transpose), static_cast<clblast::Transpose>(b_transpose), m, n, k, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, - clblast::float2{beta.s[0], beta.s[1]}, + float2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -312,10 +442,10 @@ StatusCode CLBlastZgemm(const Layout layout, const Transpose a_transpose, const static_cast<clblast::Transpose>(a_transpose), static_cast<clblast::Transpose>(b_transpose), m, n, k, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, - clblast::double2{beta.s[0], beta.s[1]}, + double2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -374,10 +504,10 @@ StatusCode CLBlastCsymm(const Layout layout, const Side side, const Triangle tri static_cast<clblast::Side>(side), static_cast<clblast::Triangle>(triangle), m, n, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, - clblast::float2{beta.s[0], beta.s[1]}, + float2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -394,10 +524,10 @@ StatusCode CLBlastZsymm(const Layout layout, const Side side, const Triangle tri static_cast<clblast::Side>(side), static_cast<clblast::Triangle>(triangle), m, n, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, - clblast::double2{beta.s[0], beta.s[1]}, + double2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -416,10 +546,10 @@ StatusCode CLBlastChemm(const Layout layout, const Side side, const Triangle tri static_cast<clblast::Side>(side), static_cast<clblast::Triangle>(triangle), m, n, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, - clblast::float2{beta.s[0], beta.s[1]}, + float2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -436,10 +566,10 @@ StatusCode CLBlastZhemm(const Layout layout, const Side side, const Triangle tri static_cast<clblast::Side>(side), static_cast<clblast::Triangle>(triangle), m, n, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, - clblast::double2{beta.s[0], beta.s[1]}, + double2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -493,9 +623,9 @@ StatusCode CLBlastCsyrk(const Layout layout, const Triangle triangle, const Tran static_cast<clblast::Triangle>(triangle), static_cast<clblast::Transpose>(a_transpose), n, k, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, - clblast::float2{beta.s[0], beta.s[1]}, + float2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -511,9 +641,9 @@ StatusCode CLBlastZsyrk(const Layout layout, const Triangle triangle, const Tran static_cast<clblast::Triangle>(triangle), static_cast<clblast::Transpose>(a_transpose), n, k, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, - clblast::double2{beta.s[0], beta.s[1]}, + double2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -610,10 +740,10 @@ StatusCode CLBlastCsyr2k(const Layout layout, const Triangle triangle, const Tra static_cast<clblast::Triangle>(triangle), static_cast<clblast::Transpose>(ab_transpose), n, k, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, - clblast::float2{beta.s[0], beta.s[1]}, + float2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -630,10 +760,10 @@ StatusCode CLBlastZsyr2k(const Layout layout, const Triangle triangle, const Tra static_cast<clblast::Triangle>(triangle), static_cast<clblast::Transpose>(ab_transpose), n, k, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, - clblast::double2{beta.s[0], beta.s[1]}, + double2{beta.s[0], beta.s[1]}, c_buffer, c_offset, c_ld, queue, event); return static_cast<StatusCode>(status); @@ -652,7 +782,7 @@ StatusCode CLBlastCher2k(const Layout layout, const Triangle triangle, const Tra static_cast<clblast::Triangle>(triangle), static_cast<clblast::Transpose>(ab_transpose), n, k, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, beta, @@ -672,7 +802,7 @@ StatusCode CLBlastZher2k(const Layout layout, const Triangle triangle, const Tra static_cast<clblast::Triangle>(triangle), static_cast<clblast::Transpose>(ab_transpose), n, k, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, beta, @@ -730,7 +860,7 @@ StatusCode CLBlastCtrmm(const Layout layout, const Side side, const Triangle tri static_cast<clblast::Transpose>(a_transpose), static_cast<clblast::Diagonal>(diagonal), m, n, - clblast::float2{alpha.s[0], alpha.s[1]}, + float2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, queue, event); @@ -748,7 +878,7 @@ StatusCode CLBlastZtrmm(const Layout layout, const Side side, const Triangle tri static_cast<clblast::Transpose>(a_transpose), static_cast<clblast::Diagonal>(diagonal), m, n, - clblast::double2{alpha.s[0], alpha.s[1]}, + double2{alpha.s[0], alpha.s[1]}, a_buffer, a_offset, a_ld, b_buffer, b_offset, b_ld, queue, event); diff --git a/src/kernels/common.opencl b/src/kernels/common.opencl index 12d63b99..8e71429e 100644 --- a/src/kernels/common.opencl +++ b/src/kernels/common.opencl @@ -109,12 +109,19 @@ R"( #define SetToOne(a) a = ONE #endif -// Multiply two complex variables (used in the define below) +// Multiply two complex variables (used in the defines below) #if PRECISION == 3232 || PRECISION == 6464 #define MulReal(a, b) a.x*b.x - a.y*b.y #define MulImag(a, b) a.x*b.y + a.y*b.x #endif +// The scalar multiply function +#if PRECISION == 3232 || PRECISION == 6464 + #define Multiply(c, a, b) c.x = MulReal(a,b); c.y = MulImag(a,b) +#else + #define Multiply(c, a, b) c = a * b +#endif + // The scalar multiply-add function #if PRECISION == 3232 || PRECISION == 6464 #define MultiplyAdd(c, a, b) c.x += MulReal(a,b); c.y += MulImag(a,b) diff --git a/src/kernels/level1/level1.opencl b/src/kernels/level1/level1.opencl index 449a20a2..7e10426b 100644 --- a/src/kernels/level1/level1.opencl +++ b/src/kernels/level1/level1.opencl @@ -46,6 +46,48 @@ R"( // ================================================================================================= +// The vectorized multiply function +inline realV MultiplyVector(realV cvec, const real aval, const realV bvec) { + #if VW == 1 + Multiply(cvec, aval, bvec); + #elif VW == 2 + Multiply(cvec.x, aval, bvec.x); + Multiply(cvec.y, aval, bvec.y); + #elif VW == 4 + Multiply(cvec.x, aval, bvec.x); + Multiply(cvec.y, aval, bvec.y); + Multiply(cvec.z, aval, bvec.z); + Multiply(cvec.w, aval, bvec.w); + #elif VW == 8 + Multiply(cvec.s0, aval, bvec.s0); + Multiply(cvec.s1, aval, bvec.s1); + Multiply(cvec.s2, aval, bvec.s2); + Multiply(cvec.s3, aval, bvec.s3); + Multiply(cvec.s4, aval, bvec.s4); + Multiply(cvec.s5, aval, bvec.s5); + Multiply(cvec.s6, aval, bvec.s6); + Multiply(cvec.s7, aval, bvec.s7); + #elif VW == 16 + Multiply(cvec.s0, aval, bvec.s0); + Multiply(cvec.s1, aval, bvec.s1); + Multiply(cvec.s2, aval, bvec.s2); + Multiply(cvec.s3, aval, bvec.s3); + Multiply(cvec.s4, aval, bvec.s4); + Multiply(cvec.s5, aval, bvec.s5); + Multiply(cvec.s6, aval, bvec.s6); + Multiply(cvec.s7, aval, bvec.s7); + Multiply(cvec.s8, aval, bvec.s8); + Multiply(cvec.s9, aval, bvec.s9); + Multiply(cvec.sA, aval, bvec.sA); + Multiply(cvec.sB, aval, bvec.sB); + Multiply(cvec.sC, aval, bvec.sC); + Multiply(cvec.sD, aval, bvec.sD); + Multiply(cvec.sE, aval, bvec.sE); + Multiply(cvec.sF, aval, bvec.sF); + #endif + return cvec; +} + // The vectorized multiply-add function inline realV MultiplyAddVector(realV cvec, const real aval, const realV bvec) { #if VW == 1 diff --git a/src/kernels/level1/xaxpy.opencl b/src/kernels/level1/xaxpy.opencl index 3d926d9e..1f1e8ce0 100644 --- a/src/kernels/level1/xaxpy.opencl +++ b/src/kernels/level1/xaxpy.opencl @@ -11,6 +11,8 @@ // strides (incx=incy=1) and no offsets (offx=offy=0). Another version is more general, but doesn't // support vector data-types. // +// This kernel uses the level-1 BLAS common tuning parameters. +// // ================================================================================================= // Enables loading of this file using the C++ pre-processor's #include (C++11 standard raw string @@ -38,8 +40,8 @@ __kernel void Xaxpy(const int n, const real alpha, // dividable by 'VW', 'WGS' and 'WPT'. __attribute__((reqd_work_group_size(WGS, 1, 1))) __kernel void XaxpyFast(const int n, const real alpha, - const __global realV* restrict xgm, - __global realV* ygm) { + const __global realV* restrict xgm, + __global realV* ygm) { #pragma unroll for (int w=0; w<WPT; ++w) { const int id = w*get_global_size(0) + get_global_id(0); diff --git a/src/kernels/level1/xcopy.opencl b/src/kernels/level1/xcopy.opencl new file mode 100644 index 00000000..97c27ccf --- /dev/null +++ b/src/kernels/level1/xcopy.opencl @@ -0,0 +1,57 @@ + +// ================================================================================================= +// 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 Xcopy kernel. It contains one fast vectorized version in case of unit +// strides (incx=incy=1) and no offsets (offx=offy=0). Another version is more general, but doesn't +// support vector data-types. +// +// This kernel uses the level-1 BLAS common tuning parameters. +// +// ================================================================================================= + +// 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"( + +// ================================================================================================= + +// Full version of the kernel with offsets and strided accesses +__attribute__((reqd_work_group_size(WGS, 1, 1))) +__kernel void Xcopy(const int n, + const __global real* restrict xgm, const int x_offset, const int x_inc, + __global real* ygm, const int y_offset, const int y_inc) { + + // Loops over the work that needs to be done (allows for an arbitrary number of threads) + #pragma unroll + for (int id = get_global_id(0); id<n; id += get_global_size(0)) { + ygm[id*y_inc + y_offset] = xgm[id*x_inc + x_offset]; + } +} + +// ================================================================================================= + +// Faster version of the kernel without offsets and strided accesses. Also assumes that 'n' is +// dividable by 'VW', 'WGS' and 'WPT'. +__attribute__((reqd_work_group_size(WGS, 1, 1))) +__kernel void XcopyFast(const int n, + const __global realV* restrict xgm, + __global realV* ygm) { + #pragma unroll + for (int w=0; w<WPT; ++w) { + const int id = w*get_global_size(0) + get_global_id(0); + ygm[id] = xgm[id]; + } +} + +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= diff --git a/src/kernels/level1/xscal.opencl b/src/kernels/level1/xscal.opencl new file mode 100644 index 00000000..956de3c0 --- /dev/null +++ b/src/kernels/level1/xscal.opencl @@ -0,0 +1,59 @@ + +// ================================================================================================= +// 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 Xscal kernel. It contains one fast vectorized version in case of unit +// strides (incx=1) and no offsets (offx=0). Another version is more general, but doesn't support +// vector data-types. +// +// This kernel uses the level-1 BLAS common tuning parameters. +// +// ================================================================================================= + +// 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"( + +// ================================================================================================= + +// Full version of the kernel with offsets and strided accesses +__attribute__((reqd_work_group_size(WGS, 1, 1))) +__kernel void Xscal(const int n, const real alpha, + __global real* xgm, const int x_offset, const int x_inc) { + + // Loops over the work that needs to be done (allows for an arbitrary number of threads) + #pragma unroll + for (int id = get_global_id(0); id<n; id += get_global_size(0)) { + real result; + Multiply(result, alpha, xgm[id*x_inc + x_offset]); + xgm[id*x_inc + x_offset] = result; + } +} + +// ================================================================================================= + +// Faster version of the kernel without offsets and strided accesses. Also assumes that 'n' is +// dividable by 'VW', 'WGS' and 'WPT'. +__attribute__((reqd_work_group_size(WGS, 1, 1))) +__kernel void XscalFast(const int n, const real alpha, + __global realV* xgm) { + #pragma unroll + for (int w=0; w<WPT; ++w) { + const int id = w*get_global_size(0) + get_global_id(0); + realV result; + result = MultiplyVector(result, alpha, xgm[id]); + xgm[id] = result; + } +} + +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= diff --git a/src/kernels/level1/xswap.opencl b/src/kernels/level1/xswap.opencl new file mode 100644 index 00000000..f6487b58 --- /dev/null +++ b/src/kernels/level1/xswap.opencl @@ -0,0 +1,61 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren <www.cedricnugteren.nl> +// +// This file contains the Xswap kernel. It contains one fast vectorized version in case of unit +// strides (incx=incy=1) and no offsets (offx=offy=0). Another version is more general, but doesn't +// support vector data-types. +// +// This kernel uses the level-1 BLAS common tuning parameters. +// +// ================================================================================================= + +// 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"( + +// ================================================================================================= + +// Full version of the kernel with offsets and strided accesses +__attribute__((reqd_work_group_size(WGS, 1, 1))) +__kernel void Xswap(const int n, + __global real* xgm, const int x_offset, const int x_inc, + __global real* ygm, const int y_offset, const int y_inc) { + + // Loops over the work that needs to be done (allows for an arbitrary number of threads) + #pragma unroll + for (int id = get_global_id(0); id<n; id += get_global_size(0)) { + real temp = xgm[id*x_inc + x_offset]; + xgm[id*x_inc + x_offset] = ygm[id*y_inc + y_offset]; + ygm[id*y_inc + y_offset] = temp; + } +} + +// ================================================================================================= + +// Faster version of the kernel without offsets and strided accesses. Also assumes that 'n' is +// dividable by 'VW', 'WGS' and 'WPT'. +__attribute__((reqd_work_group_size(WGS, 1, 1))) +__kernel void XswapFast(const int n, + __global realV* xgm, + __global realV* ygm) { + #pragma unroll + for (int w=0; w<WPT; ++w) { + const int id = w*get_global_size(0) + get_global_id(0); + realV temp = xgm[id]; + xgm[id] = ygm[id]; + ygm[id] = temp; + } +} + +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= diff --git a/src/routines/level1/xcopy.cc b/src/routines/level1/xcopy.cc new file mode 100644 index 00000000..52e029b9 --- /dev/null +++ b/src/routines/level1/xcopy.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 Xcopy class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xcopy.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xcopy<float>::precision_ = Precision::kSingle; +template <> const Precision Xcopy<double>::precision_ = Precision::kDouble; +template <> const Precision Xcopy<float2>::precision_ = Precision::kComplexSingle; +template <> const Precision Xcopy<double2>::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xcopy<T>::Xcopy(Queue &queue, Event &event): + Routine<T>(queue, event, "COPY", {"Xaxpy"}, precision_) { + source_string_ = + #include "../../kernels/level1/level1.opencl" + #include "../../kernels/level1/xcopy.opencl" + ; +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xcopy<T>::DoCopy(const size_t n, + 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) { + + // Makes sure all dimensions are larger than zero + if (n == 0) { return StatusCode::kInvalidDimension; } + + // Tests the vectors for validity + auto status = TestVectorX(n, x_buffer, x_offset, x_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestVectorY(n, y_buffer, y_offset, y_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Determines whether or not the fast-version can be used + bool use_fast_kernel = (x_offset == 0) && (x_inc == 1) && + (y_offset == 0) && (y_inc == 1) && + IsMultiple(n, db_["WGS"]*db_["WPT"]*db_["VW"]); + + // If possible, run the fast-version of the kernel + auto kernel_name = (use_fast_kernel) ? "XcopyFast" : "Xcopy"; + + // Retrieves the Xcopy kernel from the compiled binary + try { + auto& program = GetProgramFromCache(); + auto kernel = Kernel(program, kernel_name); + + // Sets the kernel arguments + if (use_fast_kernel) { + kernel.SetArgument(0, static_cast<int>(n)); + kernel.SetArgument(1, x_buffer()); + kernel.SetArgument(2, y_buffer()); + } + else { + kernel.SetArgument(0, static_cast<int>(n)); + kernel.SetArgument(1, x_buffer()); + kernel.SetArgument(2, static_cast<int>(x_offset)); + kernel.SetArgument(3, static_cast<int>(x_inc)); + kernel.SetArgument(4, y_buffer()); + kernel.SetArgument(5, static_cast<int>(y_offset)); + kernel.SetArgument(6, static_cast<int>(y_inc)); + } + + // Launches the kernel + if (use_fast_kernel) { + auto global = std::vector<size_t>{CeilDiv(n, db_["WPT"]*db_["VW"])}; + auto local = std::vector<size_t>{db_["WGS"]}; + status = RunKernel(kernel, global, local); + } + else { + auto n_ceiled = Ceil(n, db_["WGS"]*db_["WPT"]); + auto global = std::vector<size_t>{n_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 Xcopy<float>; +template class Xcopy<double>; +template class Xcopy<float2>; +template class Xcopy<double2>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/level1/xscal.cc b/src/routines/level1/xscal.cc new file mode 100644 index 00000000..13e1080c --- /dev/null +++ b/src/routines/level1/xscal.cc @@ -0,0 +1,111 @@ + +// ================================================================================================= +// 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 Xscal class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xscal.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xscal<float>::precision_ = Precision::kSingle; +template <> const Precision Xscal<double>::precision_ = Precision::kDouble; +template <> const Precision Xscal<float2>::precision_ = Precision::kComplexSingle; +template <> const Precision Xscal<double2>::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xscal<T>::Xscal(Queue &queue, Event &event): + Routine<T>(queue, event, "SCAL", {"Xaxpy"}, precision_) { + source_string_ = + #include "../../kernels/level1/level1.opencl" + #include "../../kernels/level1/xscal.opencl" + ; +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xscal<T>::DoScal(const size_t n, const T alpha, + const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc) { + + // Makes sure all dimensions are larger than zero + if (n == 0) { return StatusCode::kInvalidDimension; } + + // Tests the vector for validity + auto status = TestVectorX(n, x_buffer, x_offset, x_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Determines whether or not the fast-version can be used + bool use_fast_kernel = (x_offset == 0) && (x_inc == 1) && + IsMultiple(n, db_["WGS"]*db_["WPT"]*db_["VW"]); + + // If possible, run the fast-version of the kernel + auto kernel_name = (use_fast_kernel) ? "XscalFast" : "Xscal"; + + // Retrieves the Xscal kernel from the compiled binary + try { + auto& program = GetProgramFromCache(); + auto kernel = Kernel(program, kernel_name); + + // Sets the kernel arguments + if (use_fast_kernel) { + kernel.SetArgument(0, static_cast<int>(n)); + kernel.SetArgument(1, alpha); + kernel.SetArgument(2, x_buffer()); + } + else { + kernel.SetArgument(0, static_cast<int>(n)); + kernel.SetArgument(1, alpha); + kernel.SetArgument(2, x_buffer()); + kernel.SetArgument(3, static_cast<int>(x_offset)); + kernel.SetArgument(4, static_cast<int>(x_inc)); + } + + // Launches the kernel + if (use_fast_kernel) { + auto global = std::vector<size_t>{CeilDiv(n, db_["WPT"]*db_["VW"])}; + auto local = std::vector<size_t>{db_["WGS"]}; + status = RunKernel(kernel, global, local); + } + else { + auto n_ceiled = Ceil(n, db_["WGS"]*db_["WPT"]); + auto global = std::vector<size_t>{n_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 Xscal<float>; +template class Xscal<double>; +template class Xscal<float2>; +template class Xscal<double2>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/level1/xswap.cc b/src/routines/level1/xswap.cc new file mode 100644 index 00000000..b22b3bdb --- /dev/null +++ b/src/routines/level1/xswap.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 Xswap class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xswap.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xswap<float>::precision_ = Precision::kSingle; +template <> const Precision Xswap<double>::precision_ = Precision::kDouble; +template <> const Precision Xswap<float2>::precision_ = Precision::kComplexSingle; +template <> const Precision Xswap<double2>::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xswap<T>::Xswap(Queue &queue, Event &event): + Routine<T>(queue, event, "SWAP", {"Xaxpy"}, precision_) { + source_string_ = + #include "../../kernels/level1/level1.opencl" + #include "../../kernels/level1/xswap.opencl" + ; +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xswap<T>::DoSwap(const size_t n, + 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) { + + // Makes sure all dimensions are larger than zero + if (n == 0) { return StatusCode::kInvalidDimension; } + + // Tests the vectors for validity + auto status = TestVectorX(n, x_buffer, x_offset, x_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestVectorY(n, y_buffer, y_offset, y_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Determines whether or not the fast-version can be used + bool use_fast_kernel = (x_offset == 0) && (x_inc == 1) && + (y_offset == 0) && (y_inc == 1) && + IsMultiple(n, db_["WGS"]*db_["WPT"]*db_["VW"]); + + // If possible, run the fast-version of the kernel + auto kernel_name = (use_fast_kernel) ? "XswapFast" : "Xswap"; + + // Retrieves the Xswap kernel from the compiled binary + try { + auto& program = GetProgramFromCache(); + auto kernel = Kernel(program, kernel_name); + + // Sets the kernel arguments + if (use_fast_kernel) { + kernel.SetArgument(0, static_cast<int>(n)); + kernel.SetArgument(1, x_buffer()); + kernel.SetArgument(2, y_buffer()); + } + else { + kernel.SetArgument(0, static_cast<int>(n)); + kernel.SetArgument(1, x_buffer()); + kernel.SetArgument(2, static_cast<int>(x_offset)); + kernel.SetArgument(3, static_cast<int>(x_inc)); + kernel.SetArgument(4, y_buffer()); + kernel.SetArgument(5, static_cast<int>(y_offset)); + kernel.SetArgument(6, static_cast<int>(y_inc)); + } + + // Launches the kernel + if (use_fast_kernel) { + auto global = std::vector<size_t>{CeilDiv(n, db_["WPT"]*db_["VW"])}; + auto local = std::vector<size_t>{db_["WGS"]}; + status = RunKernel(kernel, global, local); + } + else { + auto n_ceiled = Ceil(n, db_["WGS"]*db_["WPT"]); + auto global = std::vector<size_t>{n_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 Xswap<float>; +template class Xswap<double>; +template class Xswap<float2>; +template class Xswap<double2>; + +// ================================================================================================= +} // namespace clblast diff --git a/test/correctness/routines/level1/xcopy.cc b/test/correctness/routines/level1/xcopy.cc new file mode 100644 index 00000000..8a06a722 --- /dev/null +++ b/test/correctness/routines/level1/xcopy.cc @@ -0,0 +1,32 @@ + +// ================================================================================================= +// 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 tests for the Xcopy routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/level1/xcopy.h" + +// ================================================================================================= + +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::RunTests<clblast::TestXcopy<float>, float, float>(argc, argv, false, "SCOPY"); + clblast::RunTests<clblast::TestXcopy<double>, double, double>(argc, argv, true, "DCOPY"); + clblast::RunTests<clblast::TestXcopy<float2>, float2, float2>(argc, argv, true, "CCOPY"); + clblast::RunTests<clblast::TestXcopy<double2>, double2, double2>(argc, argv, true, "ZCOPY"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/routines/level1/xscal.cc b/test/correctness/routines/level1/xscal.cc new file mode 100644 index 00000000..ceb1b7cd --- /dev/null +++ b/test/correctness/routines/level1/xscal.cc @@ -0,0 +1,32 @@ + +// ================================================================================================= +// 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 tests for the Xscal routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/level1/xscal.h" + +// ================================================================================================= + +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::RunTests<clblast::TestXscal<float>, float, float>(argc, argv, false, "SSCAL"); + clblast::RunTests<clblast::TestXscal<double>, double, double>(argc, argv, true, "DSCAL"); + clblast::RunTests<clblast::TestXscal<float2>, float2, float2>(argc, argv, true, "CSCAL"); + clblast::RunTests<clblast::TestXscal<double2>, double2, double2>(argc, argv, true, "ZSCAL"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/routines/level1/xswap.cc b/test/correctness/routines/level1/xswap.cc new file mode 100644 index 00000000..140ccf24 --- /dev/null +++ b/test/correctness/routines/level1/xswap.cc @@ -0,0 +1,32 @@ + +// ================================================================================================= +// 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 tests for the Xswap routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/level1/xswap.h" + +// ================================================================================================= + +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::RunTests<clblast::TestXswap<float>, float, float>(argc, argv, false, "SSWAP"); + clblast::RunTests<clblast::TestXswap<double>, double, double>(argc, argv, true, "DSWAP"); + clblast::RunTests<clblast::TestXswap<float2>, float2, float2>(argc, argv, true, "CSWAP"); + clblast::RunTests<clblast::TestXswap<double2>, double2, double2>(argc, argv, true, "ZSWAP"); + return 0; +} + +// ================================================================================================= diff --git a/test/performance/routines/level1/xcopy.cc b/test/performance/routines/level1/xcopy.cc new file mode 100644 index 00000000..70b6b348 --- /dev/null +++ b/test/performance/routines/level1/xcopy.cc @@ -0,0 +1,40 @@ + +// ================================================================================================= +// 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 Xcopy command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/level1/xcopy.h" + +// ================================================================================================= + +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + switch(clblast::GetPrecision(argc, argv)) { + case clblast::Precision::kHalf: + throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kSingle: + clblast::RunClient<clblast::TestXcopy<float>, float, float>(argc, argv); break; + case clblast::Precision::kDouble: + clblast::RunClient<clblast::TestXcopy<double>, double, double>(argc, argv); break; + case clblast::Precision::kComplexSingle: + clblast::RunClient<clblast::TestXcopy<float2>, float2, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXcopy<double2>, double2, double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/test/performance/routines/level1/xscal.cc b/test/performance/routines/level1/xscal.cc new file mode 100644 index 00000000..3963ba3a --- /dev/null +++ b/test/performance/routines/level1/xscal.cc @@ -0,0 +1,40 @@ + +// ================================================================================================= +// 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 Xscal command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/level1/xscal.h" + +// ================================================================================================= + +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + switch(clblast::GetPrecision(argc, argv)) { + case clblast::Precision::kHalf: + throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kSingle: + clblast::RunClient<clblast::TestXscal<float>, float, float>(argc, argv); break; + case clblast::Precision::kDouble: + clblast::RunClient<clblast::TestXscal<double>, double, double>(argc, argv); break; + case clblast::Precision::kComplexSingle: + clblast::RunClient<clblast::TestXscal<float2>, float2, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXscal<double2>, double2, double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/test/performance/routines/level1/xswap.cc b/test/performance/routines/level1/xswap.cc new file mode 100644 index 00000000..94f271ee --- /dev/null +++ b/test/performance/routines/level1/xswap.cc @@ -0,0 +1,40 @@ + +// ================================================================================================= +// 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 Xswap command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/level1/xswap.h" + +// ================================================================================================= + +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + switch(clblast::GetPrecision(argc, argv)) { + case clblast::Precision::kHalf: + throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kSingle: + clblast::RunClient<clblast::TestXswap<float>, float, float>(argc, argv); break; + case clblast::Precision::kDouble: + clblast::RunClient<clblast::TestXswap<double>, double, double>(argc, argv); break; + case clblast::Precision::kComplexSingle: + clblast::RunClient<clblast::TestXswap<float2>, float2, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXswap<double2>, double2, double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/test/routines/level1/xcopy.h b/test/routines/level1/xcopy.h new file mode 100644 index 00000000..8d324d88 --- /dev/null +++ b/test/routines/level1/xcopy.h @@ -0,0 +1,117 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren <www.cedricnugteren.nl> +// +// This file implements a class with static methods to describe the Xcopy routine. Examples of +// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These +// static methods are used by the correctness tester and the performance tester. +// +// ================================================================================================= + +#ifndef CLBLAST_TEST_ROUTINES_XCOPY_H_ +#define CLBLAST_TEST_ROUTINES_XCOPY_H_ + +#include <vector> +#include <string> + +#include "wrapper_clblas.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class TestXcopy { + public: + + // The BLAS level: 1, 2, or 3 + static size_t BLASLevel() { return 1; } + + // The list of arguments relevant for this routine + static std::vector<std::string> GetOptions() { + return {kArgN, + kArgXInc, kArgYInc, + kArgXOffset, kArgYOffset}; + } + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeX(const Arguments<T> &args) { + return args.n * args.x_inc + args.x_offset; + } + static size_t GetSizeY(const Arguments<T> &args) { + return args.n * args.y_inc + args.y_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments<T> &args) { + args.x_size = GetSizeX(args); + args.y_size = GetSizeY(args); + } + + // Describes what the default values of the leading dimensions of the matrices are + static size_t DefaultLDA(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t DefaultLDB(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine + + // Describes which transpose options are relevant for this routine + using Transposes = std::vector<Transpose>; + static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine + static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = Copy<T>(args.n, + buffers.x_vec(), args.x_offset, args.x_inc, + buffers.y_vec(), args.y_offset, args.y_inc, + &queue_plain, &event); + clWaitForEvents(1, &event); + return status; + } + + // Describes how to run the clBLAS routine (for correctness/performance comparison) + static StatusCode RunReference(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXcopy<T>(args.n, + buffers.x_vec(), args.x_offset, args.x_inc, + buffers.y_vec(), args.y_offset, args.y_inc, + 1, &queue_plain, 0, nullptr, &event); + clWaitForEvents(1, &event); + return static_cast<StatusCode>(status); + } + + // Describes how to download the results of the computation (more importantly: which buffer) + static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers<T> &buffers, Queue &queue) { + std::vector<T> result(args.y_size, static_cast<T>(0)); + buffers.y_vec.Read(queue, args.y_size, result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<T> &args) { return args.n; } + static size_t ResultID2(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t) { + return id1*args.y_inc + args.y_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<T> &args) { + return 1 * args.n; + } + static size_t GetBytes(const Arguments<T> &args) { + return (2 * args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XCOPY_H_ +#endif diff --git a/test/routines/level1/xscal.h b/test/routines/level1/xscal.h new file mode 100644 index 00000000..d990afcc --- /dev/null +++ b/test/routines/level1/xscal.h @@ -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 a class with static methods to describe the Xscal routine. Examples of +// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These +// static methods are used by the correctness tester and the performance tester. +// +// ================================================================================================= + +#ifndef CLBLAST_TEST_ROUTINES_XSCAL_H_ +#define CLBLAST_TEST_ROUTINES_XSCAL_H_ + +#include <vector> +#include <string> + +#include "wrapper_clblas.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class TestXscal { + public: + + // The BLAS level: 1, 2, or 3 + static size_t BLASLevel() { return 1; } + + // The list of arguments relevant for this routine + static std::vector<std::string> GetOptions() { + return {kArgN, + kArgXInc, + kArgXOffset, + kArgAlpha}; + } + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeX(const Arguments<T> &args) { + return args.n * args.x_inc + args.x_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments<T> &args) { + args.x_size = GetSizeX(args); + } + + // Describes what the default values of the leading dimensions of the matrices are + static size_t DefaultLDA(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t DefaultLDB(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine + + // Describes which transpose options are relevant for this routine + using Transposes = std::vector<Transpose>; + static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine + static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = Scal(args.n, args.alpha, + buffers.x_vec(), args.x_offset, args.x_inc, + &queue_plain, &event); + clWaitForEvents(1, &event); + return status; + } + + // Describes how to run the clBLAS routine (for correctness/performance comparison) + static StatusCode RunReference(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXscal(args.n, args.alpha, + buffers.x_vec(), args.x_offset, args.x_inc, + 1, &queue_plain, 0, nullptr, &event); + clWaitForEvents(1, &event); + return static_cast<StatusCode>(status); + } + + // Describes how to download the results of the computation (more importantly: which buffer) + static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers<T> &buffers, Queue &queue) { + std::vector<T> result(args.x_size, static_cast<T>(0)); + buffers.x_vec.Read(queue, args.x_size, result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<T> &args) { return args.n; } + static size_t ResultID2(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t) { + return id1*args.x_inc + args.x_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<T> &args) { + return args.n; + } + static size_t GetBytes(const Arguments<T> &args) { + return (2 * args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XSCAL_H_ +#endif diff --git a/test/routines/level1/xswap.h b/test/routines/level1/xswap.h new file mode 100644 index 00000000..2096a2c3 --- /dev/null +++ b/test/routines/level1/xswap.h @@ -0,0 +1,118 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren <www.cedricnugteren.nl> +// +// This file implements a class with static methods to describe the Xswap routine. Examples of +// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These +// static methods are used by the correctness tester and the performance tester. +// +// ================================================================================================= + +#ifndef CLBLAST_TEST_ROUTINES_XSWAP_H_ +#define CLBLAST_TEST_ROUTINES_XSWAP_H_ + +#include <vector> +#include <string> + +#include "wrapper_clblas.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class TestXswap { + public: + + // The BLAS level: 1, 2, or 3 + static size_t BLASLevel() { return 1; } + + // The list of arguments relevant for this routine + static std::vector<std::string> GetOptions() { + return {kArgN, + kArgXInc, kArgYInc, + kArgXOffset, kArgYOffset}; + } + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeX(const Arguments<T> &args) { + return args.n * args.x_inc + args.x_offset; + } + static size_t GetSizeY(const Arguments<T> &args) { + return args.n * args.y_inc + args.y_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments<T> &args) { + args.x_size = GetSizeX(args); + args.y_size = GetSizeY(args); + } + + // Describes what the default values of the leading dimensions of the matrices are + static size_t DefaultLDA(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t DefaultLDB(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine + + // Describes which transpose options are relevant for this routine + using Transposes = std::vector<Transpose>; + static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine + static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = Swap<T>(args.n, + buffers.x_vec(), args.x_offset, args.x_inc, + buffers.y_vec(), args.y_offset, args.y_inc, + &queue_plain, &event); + clWaitForEvents(1, &event); + return status; + } + + // Describes how to run the clBLAS routine (for correctness/performance comparison) + static StatusCode RunReference(const Arguments<T> &args, const Buffers<T> &buffers, Queue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXswap<T>(args.n, + buffers.x_vec(), args.x_offset, args.x_inc, + buffers.y_vec(), args.y_offset, args.y_inc, + 1, &queue_plain, 0, nullptr, &event); + clWaitForEvents(1, &event); + return static_cast<StatusCode>(status); + } + + // Describes how to download the results of the computation (more importantly: which buffer) + static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers<T> &buffers, Queue &queue) { + std::vector<T> result(args.x_size + args.y_size, static_cast<T>(0)); + buffers.x_vec.Read(queue, args.x_size, &result[0]); + buffers.y_vec.Read(queue, args.y_size, &result[args.x_size]); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<T> &args) { return args.n; } + static size_t ResultID2(const Arguments<T> &) { return 2; } // x_vec and y_vec + static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t id2) { + return (id2 == 0) ? id1*args.x_inc + args.x_offset : id1*args.y_inc + args.y_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<T> &args) { + return args.n; + } + static size_t GetBytes(const Arguments<T> &args) { + return (2 * args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XSWAP_H_ +#endif diff --git a/test/wrapper_clblas.h b/test/wrapper_clblas.h index 86810fa2..e17fae76 100644 --- a/test/wrapper_clblas.h +++ b/test/wrapper_clblas.h @@ -23,6 +23,150 @@ namespace clblast { // ================================================================================================= // BLAS level-1 (vector-vector) routines +// Calls {clblasSswap, clblasDswap, clblasCswap, clblasZswap} with the arguments forwarded. +template <typename T> clblasStatus clblasXswap( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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); +template <> clblasStatus clblasXswap<float>( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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 clblasSswap(n, + x_vec, x_offset, static_cast<int>(x_inc), + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXswap<double>( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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 clblasDswap(n, + x_vec, x_offset, static_cast<int>(x_inc), + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXswap<float2>( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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 clblasCswap(n, + x_vec, x_offset, static_cast<int>(x_inc), + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXswap<double2>( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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 clblasZswap(n, + x_vec, x_offset, static_cast<int>(x_inc), + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} + +// Calls {clblasSscal, clblasDscal, clblasCscal, clblasZscal} with the arguments forwarded. +clblasStatus clblasXscal( + size_t n, float alpha, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasSscal(n, alpha, + x_vec, x_offset, static_cast<int>(x_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXscal( + size_t n, double alpha, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasDscal(n, alpha, + x_vec, x_offset, static_cast<int>(x_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXscal( + size_t n, float2 alpha, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto cl_alpha = cl_float2{{alpha.real(), alpha.imag()}}; + return clblasCscal(n, cl_alpha, + x_vec, x_offset, static_cast<int>(x_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXscal( + size_t n, double2 alpha, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto cl_alpha = cl_double2{{alpha.real(), alpha.imag()}}; + return clblasZscal(n, cl_alpha, + x_vec, x_offset, static_cast<int>(x_inc), + num_queues, queues, num_wait_events, wait_events, events); +} + +// Calls {clblasScopy, clblasDcopy, clblasCcopy, clblasZcopy} with the arguments forwarded. +template <typename T> clblasStatus clblasXcopy( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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); +template <> clblasStatus clblasXcopy<float>( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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 clblasScopy(n, + x_vec, x_offset, static_cast<int>(x_inc), + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXcopy<double>( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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 clblasDcopy(n, + x_vec, x_offset, static_cast<int>(x_inc), + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXcopy<float2>( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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 clblasCcopy(n, + x_vec, x_offset, static_cast<int>(x_inc), + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXcopy<double2>( + size_t n, + const cl_mem x_vec, size_t x_offset, size_t x_inc, + 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 clblasZcopy(n, + x_vec, x_offset, static_cast<int>(x_inc), + y_vec, y_offset, static_cast<int>(y_inc), + num_queues, queues, num_wait_events, wait_events, events); +} + // Calls {clblasSaxpy, clblasDaxpy, clblasCaxpy, clblasZaxpy} with the arguments forwarded. clblasStatus clblasXaxpy( size_t n, float alpha, |