diff options
| author | Cedric Nugteren <web@cedricnugteren.nl> | 2015-09-14 17:12:23 +0200 |
|---|---|---|
| committer | Cedric Nugteren <web@cedricnugteren.nl> | 2015-09-14 17:12:23 +0200 |
| commit | a2b773573deddf95d9217bdbfe9223df9496c391 (patch) | |
| tree | 0a1db9d536d3d9b76139069c03032b8bc3cd3f2a | |
| parent | 70ba7c83d4b5a786264fe547f486840de594950f (diff) | |
| parent | 224c96758468b960b776debaa45445ac5288df0d (diff) | |
Merge pull request #25 from CNugteren/level1_routines
Added several level 1 routines
65 files changed, 3617 insertions, 117 deletions
@@ -1,6 +1,12 @@ Development version (next release) -- +- Added level-1 routines: + * SSWAP/DSWAP/CSWAP/ZSWAP + * SSCAL/DSCAL/CSCAL/ZSCAL + * SCOPY/DCOPY/CCOPY/ZCOPY + * SDOT/DDOT + * CDOTU/ZDOTU + * CDOTC/ZDOTC Version 0.4.0 - Now using the Claduc C++11 interface to OpenCL diff --git a/CMakeLists.txt b/CMakeLists.txt index c9a398a7..1960bf1d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -102,10 +102,10 @@ include_directories(${clblast_SOURCE_DIR}/include ${OPENCL_INCLUDE_DIRS}) # ================================================================================================== # Sets the supported routines and the used kernels. New routines and kernels should be added here. -set(KERNELS copy pad transpose padtranspose xaxpy xgemv xgemm) +set(KERNELS copy pad transpose padtranspose xaxpy xdot xgemv xgemm) set(SAMPLE_PROGRAMS_CPP sgemm) set(SAMPLE_PROGRAMS_C sgemm) -set(LEVEL1_ROUTINES xaxpy) +set(LEVEL1_ROUTINES xswap xscal xcopy xaxpy xdot xdotu xdotc) 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,14 +138,13 @@ 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 | - | - | | | | -| xDOTC | - | - | | | | -| xxxDOT | - | - | - | - | +SDS | +| xDOT | ✔ | ✔ | - | - | | +| xDOTU | - | - | ✔ | ✔ | | +| xDOTC | - | - | ✔ | ✔ | | | xNRM2 | | | - | - | +SC +DZ | | xASUM | | | - | - | +SC +DZ | | IxAMAX | | | | | | diff --git a/include/clblast.h b/include/clblast.h index bd0f161c..72825e0b 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -68,6 +68,8 @@ enum class StatusCode { kInvalidLocalMemUsage = -2046, // Not enough local memory available on this device kNoHalfPrecision = -2045, // Half precision (16-bits) not supported by the device kNoDoublePrecision = -2044, // Double precision (64-bits) not supported by the device + kInvalidVectorDot = -2043, // Vector dot is not a valid OpenCL buffer + kInsufficientMemoryDot = -2042, // Vector dot's OpenCL buffer is too small }; // Matrix layout and transpose types @@ -83,16 +85,64 @@ 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); + +// Vector copy: SCOPY/DCOPY/CCOPY/ZCOPY +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); -// Templated-precision vector-times-constant plus vector: SAXPY/DAXPY/CAXPY/ZAXPY +// Vector-times-constant plus vector: SAXPY/DAXPY/CAXPY/ZAXPY 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); +// Dot product of two vectors: SDOT/DDOT +template <typename T> +StatusCode Dot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); + +// Dot product of two complex vectors: CDOTU/ZDOTU +template <typename T> +StatusCode Dotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); + +// Dot product of two complex vectors, one conjugated: CDOTC/ZDOTC +template <typename T> +StatusCode Dotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_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 +179,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..88754990 100644 --- a/include/clblast_c.h +++ b/include/clblast_c.h @@ -64,6 +64,8 @@ typedef enum StatusCode_ { kInvalidLocalMemUsage = -2046, // Not enough local memory available on this device kNoHalfPrecision = -2045, // Half precision (16-bits) not supported by the device kNoDoublePrecision = -2044, // Double precision (64-bits) not supported by the device + kInvalidVectorDot = -2043, // Vector dot is not a valid OpenCL buffer + kInsufficientMemoryDot = -2042, // Vector dot's OpenCL buffer is too small } StatusCode; // Matrix layout and transpose types @@ -81,6 +83,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, @@ -103,6 +159,42 @@ 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); +// Dot product of two vectors: SDOT/DDOT +StatusCode CLBlastSdot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); +StatusCode CLBlastDdot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); + +// Dot product of two complex vectors: CDOTU/ZDOTU +StatusCode CLBlastCdotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); +StatusCode CLBlastZdotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); + +// Dot product of two complex vectors, one conjugated: CDOTC/ZDOTC +StatusCode CLBlastCdotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); +StatusCode CLBlastZdotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event); + // ================================================================================================= // BLAS level-2 (matrix-vector) routines // ================================================================================================= diff --git a/include/internal/database.h b/include/internal/database.h index 8c937e34..1ac0e646 100644 --- a/include/internal/database.h +++ b/include/internal/database.h @@ -67,6 +67,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 XdotSingle, XdotDouble, XdotComplexSingle, XdotComplexDouble; static const DatabaseEntry XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble; static const DatabaseEntry XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble; static const DatabaseEntry CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble; diff --git a/include/internal/database/xdot.h b/include/internal/database/xdot.h new file mode 100644 index 00000000..05841eb7 --- /dev/null +++ b/include/internal/database/xdot.h @@ -0,0 +1,113 @@ + +// ================================================================================================= +// 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 Xdot kernels. +// +// ================================================================================================= + +namespace clblast { +// ================================================================================================= + +const Database::DatabaseEntry Database::XdotSingle = { + "Xdot", Precision::kSingle, { + { // NVIDIA GPUs + kDeviceTypeGPU, kDeviceVendorNVIDIA, { + } + }, + { // AMD GPUs + kDeviceTypeGPU, kDeviceVendorAMD, { + } + }, + { // Intel GPUs + kDeviceTypeGPU, kDeviceVendorIntel, { + { "Iris", { {"WGS1",512}, {"WGS2",512} } }, + } + }, + { // Default + kDeviceTypeAll, kDeviceVendorAll, { + { kDefaultDevice, { {"WGS1",64}, {"WGS2",64} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XdotDouble = { + "Xdot", Precision::kDouble, { + { // NVIDIA GPUs + kDeviceTypeGPU, kDeviceVendorNVIDIA, { + } + }, + { // AMD GPUs + kDeviceTypeGPU, kDeviceVendorAMD, { + } + }, + { // Intel GPUs + kDeviceTypeGPU, kDeviceVendorIntel, { + } + }, + { // Default + kDeviceTypeAll, kDeviceVendorAll, { + { kDefaultDevice, { {"WGS1",64}, {"WGS2",64} } }, + } + }, + } +}; +// ================================================================================================= + +const Database::DatabaseEntry Database::XdotComplexSingle = { + "Xdot", Precision::kComplexSingle, { + { // NVIDIA GPUs + kDeviceTypeGPU, kDeviceVendorNVIDIA, { + } + }, + { // AMD GPUs + kDeviceTypeGPU, kDeviceVendorAMD, { + } + }, + { // Intel GPUs + kDeviceTypeGPU, kDeviceVendorIntel, { + { "Iris", { {"WGS1",512}, {"WGS2",512} } }, + } + }, + { // Default + kDeviceTypeAll, kDeviceVendorAll, { + { kDefaultDevice, { {"WGS1",64}, {"WGS2",64} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XdotComplexDouble = { + "Xdot", Precision::kComplexDouble, { + { // NVIDIA GPUs + kDeviceTypeGPU, kDeviceVendorNVIDIA, { + } + }, + { // AMD GPUs + kDeviceTypeGPU, kDeviceVendorAMD, { + } + }, + { // Intel GPUs + kDeviceTypeGPU, kDeviceVendorIntel, { + } + }, + { // Default + kDeviceTypeAll, kDeviceVendorAll, { + { kDefaultDevice, { {"WGS1",64}, {"WGS2",64} } }, + } + }, + } +}; + +// ================================================================================================= +} // namespace clblast diff --git a/include/internal/routine.h b/include/internal/routine.h index 367917fd..c5b253b4 100644 --- a/include/internal/routine.h +++ b/include/internal/routine.h @@ -78,6 +78,8 @@ class Routine { const size_t inc, const size_t data_size); StatusCode TestVectorY(const size_t n, const Buffer<T> &buffer, const size_t offset, const size_t inc, const size_t data_size); + StatusCode TestVectorDot(const size_t n, const Buffer<T> &buffer, const size_t offset, + const size_t data_size); // Copies/transposes a matrix and padds/unpads it with zeroes. This method is also able to write // to symmetric and triangular matrices through optional arguments. 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/xdot.h b/include/internal/routines/level1/xdot.h new file mode 100644 index 00000000..64b62945 --- /dev/null +++ b/include/internal/routines/level1/xdot.h @@ -0,0 +1,58 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren <www.cedricnugteren.nl> +// +// This file implements the Xdot routine. The precision is implemented using a template argument. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XDOT_H_ +#define CLBLAST_ROUTINES_XDOT_H_ + +#include "internal/routine.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xdot: public Routine<T> { + public: + + // Members and methods from the base class + using Routine<T>::db_; + using Routine<T>::source_string_; + using Routine<T>::queue_; + using Routine<T>::context_; + using Routine<T>::GetProgramFromCache; + using Routine<T>::TestVectorX; + using Routine<T>::TestVectorY; + using Routine<T>::TestVectorDot; + using Routine<T>::RunKernel; + using Routine<T>::ErrorIn; + + // Constructor + Xdot(Queue &queue, Event &event, const std::string &name = "DOT"); + + // Templated-precision implementation of the routine + StatusCode DoDot(const size_t n, + const Buffer<T> &dot_buffer, const size_t dot_offset, + const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc, + const Buffer<T> &y_buffer, const size_t y_offset, const size_t y_inc, + const bool do_conjugate = false); + + private: + // Static variable to get the precision + const static Precision precision_; +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XDOT_H_ +#endif diff --git a/include/internal/routines/level1/xdotc.h b/include/internal/routines/level1/xdotc.h new file mode 100644 index 00000000..726cec7c --- /dev/null +++ b/include/internal/routines/level1/xdotc.h @@ -0,0 +1,44 @@ + +// ================================================================================================= +// 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 Xdotc routine. The precision is implemented using a template argument. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XDOTC_H_ +#define CLBLAST_ROUTINES_XDOTC_H_ + +#include "internal/routines/level1/xdot.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xdotc: public Xdot<T> { + public: + + // Uses the regular Xdot routine + using Xdot<T>::DoDot; + + // Constructor + Xdotc(Queue &queue, Event &event, const std::string &name = "DOTC"); + + // Templated-precision implementation of the routine + StatusCode DoDotc(const size_t n, + const Buffer<T> &dot_buffer, const size_t dot_offset, + 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); +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XDOTC_H_ +#endif diff --git a/include/internal/routines/level1/xdotu.h b/include/internal/routines/level1/xdotu.h new file mode 100644 index 00000000..825ebb78 --- /dev/null +++ b/include/internal/routines/level1/xdotu.h @@ -0,0 +1,44 @@ + +// ================================================================================================= +// 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 Xdotu routine. The precision is implemented using a template argument. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XDOTU_H_ +#define CLBLAST_ROUTINES_XDOTU_H_ + +#include "internal/routines/level1/xdot.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xdotu: public Xdot<T> { + public: + + // Uses the regular Xdot routine + using Xdot<T>::DoDot; + + // Constructor + Xdotu(Queue &queue, Event &event, const std::string &name = "DOTU"); + + // Templated-precision implementation of the routine + StatusCode DoDotu(const size_t n, + const Buffer<T> &dot_buffer, const size_t dot_offset, + 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); +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XDOTU_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/include/internal/tuning.h b/include/internal/tuning.h index f029c704..6ea530ba 100644 --- a/include/internal/tuning.h +++ b/include/internal/tuning.h @@ -64,11 +64,13 @@ void Tuner(int argc, char* argv[]) { auto a_mat = std::vector<T>(C::GetSizeA(args)); auto b_mat = std::vector<T>(C::GetSizeB(args)); auto c_mat = std::vector<T>(C::GetSizeC(args)); + auto temp = std::vector<T>(C::GetSizeTemp(args)); PopulateVector(x_vec); PopulateVector(y_vec); PopulateVector(a_mat); PopulateVector(b_mat); PopulateVector(c_mat); + PopulateVector(temp); // Initializes the tuner for the chosen device cltune::Tuner tuner(args.platform_id, args.device_id); @@ -85,7 +87,7 @@ void Tuner(int argc, char* argv[]) { // Loads the kernel sources and defines the kernel to tune auto sources = C::GetSources(); auto id = tuner.AddKernelFromString(sources, C::KernelName(), C::GlobalSize(args), C::LocalSize()); - tuner.SetReferenceFromString(sources, C::KernelName(), C::GlobalSize(args), C::LocalSizeRef()); + tuner.SetReferenceFromString(sources, C::KernelName(), C::GlobalSizeRef(args), C::LocalSizeRef()); // Sets the tunable parameters and their possible values C::SetParameters(tuner, id); @@ -103,7 +105,7 @@ void Tuner(int argc, char* argv[]) { for (auto ¶meters: C::DivGlobal()) { tuner.DivGlobalSize(id, parameters); } // Sets the function's arguments - C::SetArguments(tuner, args, x_vec, y_vec, a_mat, b_mat, c_mat); + C::SetArguments(tuner, args, x_vec, y_vec, a_mat, b_mat, c_mat, temp); // Starts the tuning process tuner.Tune(); diff --git a/include/internal/utilities.h b/include/internal/utilities.h index d9fdb9ab..466ac810 100644 --- a/include/internal/utilities.h +++ b/include/internal/utilities.h @@ -57,6 +57,7 @@ constexpr auto kArgCLeadDim = "ldc"; constexpr auto kArgAOffset = "offa"; constexpr auto kArgBOffset = "offb"; constexpr auto kArgCOffset = "offc"; +constexpr auto kArgDotOffset = "offdot"; constexpr auto kArgAlpha = "alpha"; constexpr auto kArgBeta = "beta"; @@ -105,6 +106,7 @@ struct Arguments { size_t a_offset = 0; size_t b_offset = 0; size_t c_offset = 0; + size_t dot_offset = 0; T alpha = T{1.0}; T beta = T{1.0}; size_t x_size = 1; @@ -112,6 +114,7 @@ struct Arguments { size_t a_size = 1; size_t b_size = 1; size_t c_size = 1; + size_t dot_size = 1; // Tuner-specific arguments double fraction = 1.0; // Client-specific arguments @@ -138,6 +141,7 @@ struct Buffers { Buffer<T> a_mat; Buffer<T> b_mat; Buffer<T> c_mat; + Buffer<T> dot; }; // ================================================================================================= diff --git a/src/clblast.cc b/src/clblast.cc index 12c7b880..3303085e 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -18,7 +18,13 @@ #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" +#include "internal/routines/level1/xdot.h" +#include "internal/routines/level1/xdotu.h" +#include "internal/routines/level1/xdotc.h" // BLAS level-2 includes #include "internal/routines/level2/xgemv.h" @@ -40,41 +46,223 @@ 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); + +// DOT +template <typename T> +StatusCode Dot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xdot<T>(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoDot(n, + Buffer<T>(dot_buffer), dot_offset, + Buffer<T>(x_buffer), x_offset, x_inc, + Buffer<T>(y_buffer), y_offset, y_inc); +} +template StatusCode Dot<float>(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Dot<double>(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); + +// DOTU +template <typename T> +StatusCode Dotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xdotu<T>(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoDotu(n, + Buffer<T>(dot_buffer), dot_offset, + Buffer<T>(x_buffer), x_offset, x_inc, + Buffer<T>(y_buffer), y_offset, y_inc); +} +template StatusCode Dotu<float2>(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Dotu<double2>(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); + +// DOTC +template <typename T> +StatusCode Dotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xdotc<T>(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoDotc(n, + Buffer<T>(dot_buffer), dot_offset, + Buffer<T>(x_buffer), x_offset, x_inc, + Buffer<T>(y_buffer), y_offset, y_inc); +} +template StatusCode Dotc<float2>(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Dotc<double2>(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + 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..eccf517f 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,13 +196,91 @@ 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); return static_cast<StatusCode>(status); } +// DOT +StatusCode CLBlastSdot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dot<float>(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastDdot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dot<double>(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} + +// DOTU +StatusCode CLBlastCdotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dotu<float2>(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastZdotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dotu<double2>(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} + +// DOTC +StatusCode CLBlastCdotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dotc<float2>(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} +StatusCode CLBlastZdotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dotc<double2>(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast<StatusCode>(status); +} + // ================================================================================================= // BLAS level-2 (matrix-vector) routines // ================================================================================================= @@ -127,10 +335,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 +354,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 +375,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 +394,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 +500,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 +520,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 +582,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 +602,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 +624,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 +644,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 +701,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 +719,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 +818,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 +838,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 +860,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 +880,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 +938,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 +956,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/database.cc b/src/database.cc index 258d861e..b7275dad 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/xdot.h" #include "internal/database/xgemv.h" #include "internal/database/xgemm.h" #include "internal/database/copy.h" @@ -28,6 +29,7 @@ namespace clblast { // Initializes the database const std::vector<Database::DatabaseEntry> Database::database = { XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble, + XdotSingle, XdotDouble, XdotComplexSingle, XdotComplexDouble, XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble, XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble, CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble, diff --git a/src/kernels/common.opencl b/src/kernels/common.opencl index 12d63b99..f2a2e7a7 100644 --- a/src/kernels/common.opencl +++ b/src/kernels/common.opencl @@ -109,12 +109,26 @@ R"( #define SetToOne(a) a = ONE #endif -// Multiply two complex variables (used in the define below) +// Adds two complex variables +#if PRECISION == 3232 || PRECISION == 6464 + #define Add(c, a, b) c.x = a.x + b.x; c.y = a.y + b.y +#else + #define Add(c, a, b) c = a + b +#endif + +// 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/xaxpy.opencl b/src/kernels/level1/level1.opencl index b7ffe9ff..7e10426b 100644 --- a/src/kernels/xaxpy.opencl +++ b/src/kernels/level1/level1.opencl @@ -7,9 +7,7 @@ // Author(s): // Cedric Nugteren <www.cedricnugteren.nl> // -// This file contains the Xaxpy 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 file contains the common functions and parameters specific for level 1 BLAS kernels. // // ================================================================================================= @@ -48,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 @@ -92,36 +132,6 @@ inline realV MultiplyAddVector(realV cvec, const real aval, const realV bvec) { // ================================================================================================= -// Full version of the kernel with offsets and strided accesses -__attribute__((reqd_work_group_size(WGS, 1, 1))) -__kernel void Xaxpy(const int n, const real alpha, - 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)) { - MultiplyAdd(ygm[id*y_inc + y_offset], alpha, 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 XaxpyFast(const int n, const real alpha, - 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] = MultiplyAddVector(ygm[id], alpha, xgm[id]); - } -} - -// ================================================================================================= - // End of the C++11 raw string literal )" diff --git a/src/kernels/level1/xaxpy.opencl b/src/kernels/level1/xaxpy.opencl new file mode 100644 index 00000000..1f1e8ce0 --- /dev/null +++ b/src/kernels/level1/xaxpy.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 Xaxpy 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 Xaxpy(const int n, const real alpha, + 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)) { + MultiplyAdd(ygm[id*y_inc + y_offset], alpha, 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 XaxpyFast(const int n, const real alpha, + 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] = MultiplyAddVector(ygm[id], alpha, xgm[id]); + } +} + +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= 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/xdot.opencl b/src/kernels/level1/xdot.opencl new file mode 100644 index 00000000..e13eb3c1 --- /dev/null +++ b/src/kernels/level1/xdot.opencl @@ -0,0 +1,106 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren <www.cedricnugteren.nl> +// +// This file contains the Xdot kernel. It implements a dot-product computation using reduction +// kernels. Reduction is split in two parts. In the first (main) kernel the X and Y vectors are +// multiplied, followed by a per-thread and a per-workgroup reduction. The second (epilogue) kernel +// is executed with a single workgroup only, computing the final result. +// +// ================================================================================================= + +// Enables loading of this file using the C++ pre-processor's #include (C++11 standard raw string +// literal). Comment-out this line for syntax-highlighting when developing. +R"( + +// Parameters set by the tuner or by the database. Here they are given a basic default value in case +// this kernel file is used outside of the CLBlast library. +#ifndef WGS1 + #define WGS1 64 // The local work-group size of the main kernel +#endif +#ifndef WGS2 + #define WGS2 64 // The local work-group size of the epilogue kernel +#endif + +// ================================================================================================= + +// The main reduction kernel, performing the multiplication and the majority of the sum operation +__attribute__((reqd_work_group_size(WGS1, 1, 1))) +__kernel void Xdot(const int n, + const __global real* restrict xgm, const int x_offset, const int x_inc, + const __global real* restrict ygm, const int y_offset, const int y_inc, + __global real* output, const int do_conjugate) { + __local real lm[WGS1]; + const int lid = get_local_id(0); + const int wgid = get_group_id(0); + const int num_groups = get_num_groups(0); + + // Performs multiplication and the first steps of the reduction + real acc; + SetToZero(acc); + int id = wgid*WGS1 + lid; + while (id < n) { + real x = xgm[id*x_inc + x_offset]; + real y = ygm[id*y_inc + y_offset]; + if (do_conjugate) { COMPLEX_CONJUGATE(x); } + MultiplyAdd(acc, x, y); + id += WGS1*num_groups; + } + lm[lid] = acc; + barrier(CLK_LOCAL_MEM_FENCE); + + // Performs reduction in local memory + #pragma unroll + for (int s=WGS1/2; s>0; s=s>>1) { + if (lid < s) { + Add(lm[lid], lm[lid], lm[lid + s]); + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + // Stores the per-workgroup result + if (lid == 0) { + output[wgid] = lm[0]; + } +} + +// ================================================================================================= + +// The epilogue reduction kernel, performing the final bit of the sum operation. This kernel has to +// be launched with a single workgroup only. +__attribute__((reqd_work_group_size(WGS2, 1, 1))) +__kernel void XdotEpilogue(const __global real* restrict input, + __global real* dot, const int dot_offset) { + __local real lm[WGS2]; + const int lid = get_local_id(0); + + // Performs the first step of the reduction while loading the data + Add(lm[lid], input[lid], input[lid + WGS2]); + barrier(CLK_LOCAL_MEM_FENCE); + + // Performs reduction in local memory + #pragma unroll + for (int s=WGS2/2; s>0; s=s>>1) { + if (lid < s) { + Add(lm[lid], lm[lid], lm[lid + s]); + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + // Stores the final result + if (lid == 0) { + dot[dot_offset] = lm[0]; + } +} + +// ================================================================================================= + +// 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/routine.cc b/src/routine.cc index 31476c42..05a03683 100644 --- a/src/routine.cc +++ b/src/routine.cc @@ -223,6 +223,21 @@ StatusCode Routine<T>::TestVectorY(const size_t n, const Buffer<T> &buffer, cons // ================================================================================================= +// Tests vector dot for validity: checks for a valid increment, a valid OpenCL buffer, and for a +// sufficient buffer size. +template <typename T> +StatusCode Routine<T>::TestVectorDot(const size_t n, const Buffer<T> &buffer, const size_t offset, + const size_t data_size) { + try { + auto required_size = (n + offset)*data_size; + auto buffer_size = buffer.GetSize(); + if (buffer_size < required_size) { return StatusCode::kInsufficientMemoryDot; } + } catch (...) { return StatusCode::kInvalidVectorDot; } + return StatusCode::kSuccess; +} + +// ================================================================================================= + // Copies or transposes a matrix and pads/unpads it with zeros template <typename T> StatusCode Routine<T>::PadCopyTransposeMatrix(const size_t src_one, const size_t src_two, diff --git a/src/routines/level1/xaxpy.cc b/src/routines/level1/xaxpy.cc index 7646b0e4..ce138fa6 100644 --- a/src/routines/level1/xaxpy.cc +++ b/src/routines/level1/xaxpy.cc @@ -32,7 +32,8 @@ template <typename T> Xaxpy<T>::Xaxpy(Queue &queue, Event &event): Routine<T>(queue, event, "AXPY", {"Xaxpy"}, precision_) { source_string_ = - #include "../../kernels/xaxpy.opencl" + #include "../../kernels/level1/level1.opencl" + #include "../../kernels/level1/xaxpy.opencl" ; } 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/xdot.cc b/src/routines/level1/xdot.cc new file mode 100644 index 00000000..a0c1e756 --- /dev/null +++ b/src/routines/level1/xdot.cc @@ -0,0 +1,115 @@ + +// ================================================================================================= +// 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 Xdot class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xdot.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xdot<float>::precision_ = Precision::kSingle; +template <> const Precision Xdot<double>::precision_ = Precision::kDouble; +template <> const Precision Xdot<float2>::precision_ = Precision::kComplexSingle; +template <> const Precision Xdot<double2>::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xdot<T>::Xdot(Queue &queue, Event &event, const std::string &name): + Routine<T>(queue, event, name, {"Xdot"}, precision_) { + source_string_ = + #include "../../kernels/level1/xdot.opencl" + ; +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xdot<T>::DoDot(const size_t n, + const Buffer<T> &dot_buffer, const size_t dot_offset, + const Buffer<T> &x_buffer, const size_t x_offset, const size_t x_inc, + const Buffer<T> &y_buffer, const size_t y_offset, const size_t y_inc, + const bool do_conjugate) { + + // 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; } + status = TestVectorDot(1, dot_buffer, dot_offset, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Retrieves the Xdot kernels from the compiled binary + try { + auto& program = GetProgramFromCache(); + auto kernel1 = Kernel(program, "Xdot"); + auto kernel2 = Kernel(program, "XdotEpilogue"); + + // Creates the buffer for intermediate values + auto temp_size = 2*db_["WGS2"]; + auto temp_buffer = Buffer<T>(context_, temp_size); + + // Sets the kernel arguments + kernel1.SetArgument(0, static_cast<int>(n)); + kernel1.SetArgument(1, x_buffer()); + kernel1.SetArgument(2, static_cast<int>(x_offset)); + kernel1.SetArgument(3, static_cast<int>(x_inc)); + kernel1.SetArgument(4, y_buffer()); + kernel1.SetArgument(5, static_cast<int>(y_offset)); + kernel1.SetArgument(6, static_cast<int>(y_inc)); + kernel1.SetArgument(7, temp_buffer()); + kernel1.SetArgument(8, static_cast<int>(do_conjugate)); + + // Launches the main kernel + auto global1 = std::vector<size_t>{db_["WGS1"]*temp_size}; + auto local1 = std::vector<size_t>{db_["WGS1"]}; + status = RunKernel(kernel1, global1, local1); + if (ErrorIn(status)) { return status; } + + // Sets the arguments for the epilogue kernel + kernel2.SetArgument(0, temp_buffer()); + kernel2.SetArgument(1, dot_buffer()); + kernel2.SetArgument(2, static_cast<int>(dot_offset)); + + // Launches the epilogue kernel + auto global2 = std::vector<size_t>{db_["WGS2"]}; + auto local2 = std::vector<size_t>{db_["WGS2"]}; + status = RunKernel(kernel2, global2, local2); + if (ErrorIn(status)) { return status; } + + // Waits for all kernels to finish + queue_.Finish(); + + // Succesfully finished the computation + return StatusCode::kSuccess; + } catch (...) { return StatusCode::kInvalidKernel; } +} + +// ================================================================================================= + +// Compiles the templated class +template class Xdot<float>; +template class Xdot<double>; +template class Xdot<float2>; +template class Xdot<double2>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/level1/xdotc.cc b/src/routines/level1/xdotc.cc new file mode 100644 index 00000000..f414f556 --- /dev/null +++ b/src/routines/level1/xdotc.cc @@ -0,0 +1,49 @@ + +// ================================================================================================= +// 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 Xdotc class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xdotc.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xdotc<T>::Xdotc(Queue &queue, Event &event, const std::string &name): + Xdot<T>(queue, event, name) { +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xdotc<T>::DoDotc(const size_t n, + const Buffer<T> &dot_buffer, const size_t dot_offset, + 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) { + return DoDot(n, dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + true); +} + +// ================================================================================================= + +// Compiles the templated class +template class Xdotc<float2>; +template class Xdotc<double2>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/level1/xdotu.cc b/src/routines/level1/xdotu.cc new file mode 100644 index 00000000..0b1bd2a8 --- /dev/null +++ b/src/routines/level1/xdotu.cc @@ -0,0 +1,49 @@ + +// ================================================================================================= +// 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 Xdotu class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xdotu.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xdotu<T>::Xdotu(Queue &queue, Event &event, const std::string &name): + Xdot<T>(queue, event, name) { +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xdotu<T>::DoDotu(const size_t n, + const Buffer<T> &dot_buffer, const size_t dot_offset, + 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) { + return DoDot(n, dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + false); +} + +// ================================================================================================= + +// Compiles the templated class +template class Xdotu<float2>; +template class Xdotu<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/src/tuning/copy.cc b/src/tuning/copy.cc index f38a28f3..23828b25 100644 --- a/src/tuning/copy.cc +++ b/src/tuning/copy.cc @@ -53,6 +53,7 @@ class TuneCopy { static size_t GetSizeA(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeB(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeC(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeTemp(const Arguments<T> &) { return 1; } // N/A for this kernel // Sets the tuning parameters and their possible values static void SetParameters(cltune::Tuner &tuner, const size_t id) { @@ -68,6 +69,7 @@ class TuneCopy { // Sets the base thread configuration static std::vector<size_t> GlobalSize(const Arguments<T> &args) { return {args.m, args.n}; } + static std::vector<size_t> GlobalSizeRef(const Arguments<T> &args) { return GlobalSize(args); } static std::vector<size_t> LocalSize() { return {1, 1}; } static std::vector<size_t> LocalSizeRef() { return {8, 8}; } @@ -81,7 +83,8 @@ class TuneCopy { // Sets the kernel's arguments static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args, std::vector<T> &, std::vector<T> &, - std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &) { + std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &, + std::vector<T> &) { tuner.AddArgumentScalar(static_cast<int>(args.m)); tuner.AddArgumentInput(a_mat); tuner.AddArgumentOutput(b_mat); diff --git a/src/tuning/pad.cc b/src/tuning/pad.cc index 2ce566fb..6a826b6b 100644 --- a/src/tuning/pad.cc +++ b/src/tuning/pad.cc @@ -53,6 +53,7 @@ class TunePad { static size_t GetSizeA(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeB(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeC(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeTemp(const Arguments<T> &) { return 1; } // N/A for this kernel // Sets the tuning parameters and their possible values static void SetParameters(cltune::Tuner &tuner, const size_t id) { @@ -68,6 +69,7 @@ class TunePad { // Sets the base thread configuration static std::vector<size_t> GlobalSize(const Arguments<T> &args) { return {args.m, args.n}; } + static std::vector<size_t> GlobalSizeRef(const Arguments<T> &args) { return GlobalSize(args); } static std::vector<size_t> LocalSize() { return {1, 1}; } static std::vector<size_t> LocalSizeRef() { return {8, 8}; } @@ -81,7 +83,8 @@ class TunePad { // Sets the kernel's arguments static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args, std::vector<T> &, std::vector<T> &, - std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &) { + std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &, + std::vector<T> &) { tuner.AddArgumentScalar(static_cast<int>(args.m)); tuner.AddArgumentScalar(static_cast<int>(args.n)); tuner.AddArgumentScalar(static_cast<int>(args.m)); diff --git a/src/tuning/padtranspose.cc b/src/tuning/padtranspose.cc index 8d494745..3f233809 100644 --- a/src/tuning/padtranspose.cc +++ b/src/tuning/padtranspose.cc @@ -53,6 +53,7 @@ class TunePadTranspose { static size_t GetSizeA(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeB(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeC(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeTemp(const Arguments<T> &) { return 1; } // N/A for this kernel // Sets the tuning parameters and their possible values static void SetParameters(cltune::Tuner &tuner, const size_t id) { @@ -72,6 +73,7 @@ class TunePadTranspose { // Sets the base thread configuration static std::vector<size_t> GlobalSize(const Arguments<T> &args) { return {args.m, args.n}; } + static std::vector<size_t> GlobalSizeRef(const Arguments<T> &args) { return GlobalSize(args); } static std::vector<size_t> LocalSize() { return {1, 1}; } static std::vector<size_t> LocalSizeRef() { return {8, 8}; } @@ -85,7 +87,8 @@ class TunePadTranspose { // Sets the kernel's arguments static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args, std::vector<T> &, std::vector<T> &, - std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &) { + std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &, + std::vector<T> &) { tuner.AddArgumentScalar(static_cast<int>(args.m)); tuner.AddArgumentScalar(static_cast<int>(args.n)); tuner.AddArgumentScalar(static_cast<int>(args.m)); diff --git a/src/tuning/transpose.cc b/src/tuning/transpose.cc index 2ffdb7aa..3998ba66 100644 --- a/src/tuning/transpose.cc +++ b/src/tuning/transpose.cc @@ -53,6 +53,7 @@ class TuneTranspose { static size_t GetSizeA(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeB(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeC(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeTemp(const Arguments<T> &) { return 1; } // N/A for this kernel // Sets the tuning parameters and their possible values static void SetParameters(cltune::Tuner &tuner, const size_t id) { @@ -73,6 +74,7 @@ class TuneTranspose { // Sets the base thread configuration static std::vector<size_t> GlobalSize(const Arguments<T> &args) { return {args.m, args.n}; } + static std::vector<size_t> GlobalSizeRef(const Arguments<T> &args) { return GlobalSize(args); } static std::vector<size_t> LocalSize() { return {1, 1}; } static std::vector<size_t> LocalSizeRef() { return {8, 8}; } @@ -86,7 +88,8 @@ class TuneTranspose { // Sets the kernel's arguments static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args, std::vector<T> &, std::vector<T> &, - std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &) { + std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &, + std::vector<T> &) { tuner.AddArgumentScalar(static_cast<int>(args.m)); tuner.AddArgumentInput(a_mat); tuner.AddArgumentOutput(b_mat); diff --git a/src/tuning/xaxpy.cc b/src/tuning/xaxpy.cc index cc9e81d3..31aa6a8e 100644 --- a/src/tuning/xaxpy.cc +++ b/src/tuning/xaxpy.cc @@ -31,7 +31,8 @@ class TuneXaxpy { static std::string GetSources() { return #include "../src/kernels/common.opencl" - #include "../src/kernels/xaxpy.opencl" + #include "../src/kernels/level1/level1.opencl" + #include "../src/kernels/level1/xaxpy.opencl" ; } @@ -52,11 +53,12 @@ class TuneXaxpy { static double DefaultFraction() { return 1.0; } // N/A for this kernel // Describes how to obtain the sizes of the buffers - static size_t GetSizeX(const Arguments<T> &args) { return args.n; } // N/A for this kernel - static size_t GetSizeY(const Arguments<T> &args) { return args.n; } // N/A for this kernel + static size_t GetSizeX(const Arguments<T> &args) { return args.n; } + static size_t GetSizeY(const Arguments<T> &args) { return args.n; } static size_t GetSizeA(const Arguments<T> &) { return 1; } // N/A for this kernel static size_t GetSizeB(const Arguments<T> &) { return 1; } // N/A for this kernel static size_t GetSizeC(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeTemp(const Arguments<T> &) { return 1; } // N/A for this kernel // Sets the tuning parameters and their possible values static void SetParameters(cltune::Tuner &tuner, const size_t id) { @@ -71,6 +73,7 @@ class TuneXaxpy { // Sets the base thread configuration static std::vector<size_t> GlobalSize(const Arguments<T> &args) { return {args.n}; } + static std::vector<size_t> GlobalSizeRef(const Arguments<T> &args) { return GlobalSize(args); } static std::vector<size_t> LocalSize() { return {1}; } static std::vector<size_t> LocalSizeRef() { return {64}; } @@ -84,7 +87,8 @@ class TuneXaxpy { // Sets the kernel's arguments static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args, std::vector<T> &x_vec, std::vector<T> &y_vec, - std::vector<T> &, std::vector<T> &, std::vector<T> &) { + std::vector<T> &, std::vector<T> &, std::vector<T> &, + std::vector<T> &) { tuner.AddArgumentScalar(static_cast<int>(args.n)); tuner.AddArgumentScalar(args.alpha); tuner.AddArgumentInput(x_vec); diff --git a/src/tuning/xdot.cc b/src/tuning/xdot.cc new file mode 100644 index 00000000..ff6bee16 --- /dev/null +++ b/src/tuning/xdot.cc @@ -0,0 +1,125 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren <www.cedricnugteren.nl> +// +// This file uses the CLTune auto-tuner to tune the xdot OpenCL kernels. Note that the results are +// not verified, since the result is not final and depends on the WGS2 parameter. +// +// ================================================================================================= + +#include <string> +#include <vector> + +#include "internal/utilities.h" +#include "internal/tuning.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class TuneXdot { + public: + + // The representative kernel and the source code + static std::string KernelFamily() { return "xdot"; } + static std::string KernelName() { return "Xdot"; } + static std::string GetSources() { + return + #include "../src/kernels/common.opencl" + #include "../src/kernels/level1/xdot.opencl" + ; + } + + // The list of arguments relevant for this routine + static std::vector<std::string> GetOptions() { return {kArgN}; } + + // Tests for valid arguments + static void TestValidArguments(const Arguments<T> &) { } + + // Sets the default values for the arguments + static size_t DefaultM() { return 1; } // N/A for this kernel + static size_t DefaultN() { return 4096*1024; } + static size_t DefaultK() { return 1; } // N/A for this kernel + static double DefaultFraction() { return 1.0; } // N/A for this kernel + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeX(const Arguments<T> &args) { return args.n; } + static size_t GetSizeY(const Arguments<T> &args) { return args.n; } + static size_t GetSizeA(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeB(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeC(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeTemp(const Arguments<T> &args) { return args.n; } // Worst case + + // Sets the tuning parameters and their possible values + static void SetParameters(cltune::Tuner &tuner, const size_t id) { + tuner.AddParameter(id, "WGS1", {32, 64, 128, 256, 512, 1024}); + tuner.AddParameter(id, "WGS2", {32, 64, 128, 256, 512, 1024}); + tuner.AddParameter(id, "VW", {1}); + } + + // Sets the constraints and local memory size + static void SetConstraints(cltune::Tuner &, const size_t) { } + static void SetLocalMemorySize(cltune::Tuner &, const size_t, const Arguments<T> &) { } + + // Sets the base thread configuration + static std::vector<size_t> GlobalSize(const Arguments<T> &) { return {2}; } + static std::vector<size_t> GlobalSizeRef(const Arguments<T> &) { return {2*64*64}; } + static std::vector<size_t> LocalSize() { return {1}; } + static std::vector<size_t> LocalSizeRef() { return {64}; } + + // Transforms the thread configuration based on the parameters + using TransformVector = std::vector<std::vector<std::string>>; + static TransformVector MulLocal() { return {{"WGS1"}}; } + static TransformVector DivLocal() { return {}; } + static TransformVector MulGlobal() { return {{"WGS1"},{"WGS2"}}; } + static TransformVector DivGlobal() { return {}; } + + // Sets the kernel's arguments + static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args, + std::vector<T> &x_vec, std::vector<T> &y_vec, + std::vector<T> &, std::vector<T> &, std::vector<T> &, + std::vector<T> &temp) { + tuner.AddArgumentScalar(static_cast<int>(args.n)); + tuner.AddArgumentInput(x_vec); + tuner.AddArgumentScalar(0); + tuner.AddArgumentScalar(1); + tuner.AddArgumentInput(y_vec); + tuner.AddArgumentScalar(0); + tuner.AddArgumentScalar(1); + tuner.AddArgumentInput(temp); // No output checking for the result - size varies + tuner.AddArgumentScalar(static_cast<int>(false)); + } + + // Describes how to compute the performance metrics + static size_t GetMetric(const Arguments<T> &args) { + return (2*args.n + 1) * GetBytes(args.precision); + } + static std::string PerformanceUnit() { return "GB/s"; } +}; + +// ================================================================================================= +} // namespace clblast + +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + switch(clblast::GetPrecision(argc, argv)) { + case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kSingle: clblast::Tuner<clblast::TuneXdot<float>, float>(argc, argv); break; + case clblast::Precision::kDouble: clblast::Tuner<clblast::TuneXdot<double>, double>(argc, argv); break; + case clblast::Precision::kComplexSingle: clblast::Tuner<clblast::TuneXdot<float2>, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: clblast::Tuner<clblast::TuneXdot<double2>, double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/src/tuning/xgemm.cc b/src/tuning/xgemm.cc index 302f2bd5..e820cfb0 100644 --- a/src/tuning/xgemm.cc +++ b/src/tuning/xgemm.cc @@ -55,6 +55,7 @@ class TuneXgemm { static size_t GetSizeA(const Arguments<T> &args) { return args.m * args.k; } static size_t GetSizeB(const Arguments<T> &args) { return args.n * args.k; } static size_t GetSizeC(const Arguments<T> &args) { return args.m * args.n; } + static size_t GetSizeTemp(const Arguments<T> &) { return 1; } // N/A for this kernel // Sets the tuning parameters and their possible values static void SetParameters(cltune::Tuner &tuner, const size_t id) { @@ -103,6 +104,7 @@ class TuneXgemm { // Sets the base thread configuration static std::vector<size_t> GlobalSize(const Arguments<T> &args) { return {args.m, args.n}; } + static std::vector<size_t> GlobalSizeRef(const Arguments<T> &args) { return GlobalSize(args); } static std::vector<size_t> LocalSize() { return {1, 1}; } static std::vector<size_t> LocalSizeRef() { return {8, 8}; } @@ -116,7 +118,8 @@ class TuneXgemm { // Sets the kernel's arguments static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args, std::vector<T> &, std::vector<T> &, - std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &c_mat) { + std::vector<T> &a_mat, std::vector<T> &b_mat, std::vector<T> &c_mat, + std::vector<T> &) { tuner.AddArgumentScalar(static_cast<int>(args.m)); tuner.AddArgumentScalar(static_cast<int>(args.n)); tuner.AddArgumentScalar(static_cast<int>(args.k)); diff --git a/src/tuning/xgemv.cc b/src/tuning/xgemv.cc index e22b5103..3d6fe595 100644 --- a/src/tuning/xgemv.cc +++ b/src/tuning/xgemv.cc @@ -56,6 +56,7 @@ class TuneXgemv { static size_t GetSizeA(const Arguments<T> &args) { return args.m * args.n; } static size_t GetSizeB(const Arguments<T> &) { return 1; } // N/A for this kernel static size_t GetSizeC(const Arguments<T> &) { return 1; } // N/A for this kernel + static size_t GetSizeTemp(const Arguments<T> &) { return 1; } // N/A for this kernel // Sets the tuning parameters and their possible values static void SetParameters(cltune::Tuner &tuner, const size_t id) { @@ -75,6 +76,7 @@ class TuneXgemv { // Sets the base thread configuration static std::vector<size_t> GlobalSize(const Arguments<T> &args) { return {args.m}; } + static std::vector<size_t> GlobalSizeRef(const Arguments<T> &args) { return GlobalSize(args); } static std::vector<size_t> LocalSize() { return {1}; } static std::vector<size_t> LocalSizeRef() { return {64}; } @@ -88,7 +90,8 @@ class TuneXgemv { // Sets the kernel's arguments static void SetArguments(cltune::Tuner &tuner, const Arguments<T> &args, std::vector<T> &x_vec, std::vector<T> &y_vec, - std::vector<T> &a_mat, std::vector<T> &, std::vector<T> &) { + std::vector<T> &a_mat, std::vector<T> &, std::vector<T> &, + std::vector<T> &) { auto a_rotated = (V==3) ? 1 : 0; tuner.AddArgumentScalar(static_cast<int>(args.m)); tuner.AddArgumentScalar(static_cast<int>(args.n)); 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/xdot.cc b/test/correctness/routines/level1/xdot.cc new file mode 100644 index 00000000..e1b70cb2 --- /dev/null +++ b/test/correctness/routines/level1/xdot.cc @@ -0,0 +1,30 @@ + +// ================================================================================================= +// 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 Xdot routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/level1/xdot.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::TestXdot<float>, float, float>(argc, argv, false, "SDOT"); + clblast::RunTests<clblast::TestXdot<double>, double, double>(argc, argv, true, "DDOT"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/routines/level1/xdotc.cc b/test/correctness/routines/level1/xdotc.cc new file mode 100644 index 00000000..15a2f88b --- /dev/null +++ b/test/correctness/routines/level1/xdotc.cc @@ -0,0 +1,30 @@ + +// ================================================================================================= +// 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 Xdotc routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/level1/xdotc.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::TestXdotc<float2>, float2, float2>(argc, argv, true, "CDOTC"); + clblast::RunTests<clblast::TestXdotc<double2>, double2, double2>(argc, argv, true, "ZDOTC"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/routines/level1/xdotu.cc b/test/correctness/routines/level1/xdotu.cc new file mode 100644 index 00000000..c8af0388 --- /dev/null +++ b/test/correctness/routines/level1/xdotu.cc @@ -0,0 +1,30 @@ + +// ================================================================================================= +// 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 Xdotu routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/level1/xdotu.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::TestXdotu<float2>, float2, float2>(argc, argv, true, "CDOTU"); + clblast::RunTests<clblast::TestXdotu<double2>, double2, double2>(argc, argv, true, "ZDOTU"); + 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/correctness/testblas.cc b/test/correctness/testblas.cc index ff81f4c3..839ac3b1 100644 --- a/test/correctness/testblas.cc +++ b/test/correctness/testblas.cc @@ -57,11 +57,13 @@ TestBlas<T,U>::TestBlas(int argc, char *argv[], const bool silent, a_source_.resize(std::max(max_mat, max_matvec)*std::max(max_ld, max_matvec) + max_offset); b_source_.resize(std::max(max_mat, max_matvec)*std::max(max_ld, max_matvec) + max_offset); c_source_.resize(std::max(max_mat, max_matvec)*std::max(max_ld, max_matvec) + max_offset); + dot_source_.resize(std::max(max_mat, max_matvec) + max_offset); PopulateVector(x_source_); PopulateVector(y_source_); PopulateVector(a_source_); PopulateVector(b_source_); PopulateVector(c_source_); + PopulateVector(dot_source_); } // =============================================================================================== @@ -81,12 +83,14 @@ void TestBlas<T,U>::TestRegular(std::vector<Arguments<U>> &test_vector, const st auto a_mat1 = Buffer<T>(context_, args.a_size); auto b_mat1 = Buffer<T>(context_, args.b_size); auto c_mat1 = Buffer<T>(context_, args.c_size); + auto dot1 = Buffer<T>(context_, args.dot_size); x_vec1.Write(queue_, args.x_size, x_source_); y_vec1.Write(queue_, args.y_size, y_source_); a_mat1.Write(queue_, args.a_size, a_source_); b_mat1.Write(queue_, args.b_size, b_source_); c_mat1.Write(queue_, args.c_size, c_source_); - auto buffers1 = Buffers<T>{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1}; + dot1.Write(queue_, args.dot_size, dot_source_); + auto buffers1 = Buffers<T>{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, dot1}; auto status1 = run_reference_(args, buffers1, queue_); // Runs the CLBlast code @@ -95,12 +99,14 @@ void TestBlas<T,U>::TestRegular(std::vector<Arguments<U>> &test_vector, const st auto a_mat2 = Buffer<T>(context_, args.a_size); auto b_mat2 = Buffer<T>(context_, args.b_size); auto c_mat2 = Buffer<T>(context_, args.c_size); + auto dot2 = Buffer<T>(context_, args.dot_size); x_vec2.Write(queue_, args.x_size, x_source_); y_vec2.Write(queue_, args.y_size, y_source_); a_mat2.Write(queue_, args.a_size, a_source_); b_mat2.Write(queue_, args.b_size, b_source_); c_mat2.Write(queue_, args.c_size, c_source_); - auto buffers2 = Buffers<T>{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2}; + dot2.Write(queue_, args.dot_size, dot_source_); + auto buffers2 = Buffers<T>{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, dot2}; auto status2 = run_routine_(args, buffers2, queue_); // Tests for equality of the two status codes @@ -149,25 +155,31 @@ void TestBlas<T,U>::TestInvalid(std::vector<Arguments<U>> &test_vector, const st auto a1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.a_size*sizeof(T), nullptr,nullptr); auto b1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.b_size*sizeof(T), nullptr,nullptr); auto c1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.c_size*sizeof(T), nullptr,nullptr); + auto d1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.dot_size*sizeof(T), nullptr,nullptr); auto x_vec1 = Buffer<T>(x1); auto y_vec1 = Buffer<T>(y1); auto a_mat1 = Buffer<T>(a1); auto b_mat1 = Buffer<T>(b1); auto c_mat1 = Buffer<T>(c1); + auto dot1 = Buffer<T>(d1); auto x2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.x_size*sizeof(T), nullptr,nullptr); auto y2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.y_size*sizeof(T), nullptr,nullptr); auto a2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.a_size*sizeof(T), nullptr,nullptr); auto b2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.b_size*sizeof(T), nullptr,nullptr); auto c2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.c_size*sizeof(T), nullptr,nullptr); + auto d2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.dot_size*sizeof(T), nullptr,nullptr); auto x_vec2 = Buffer<T>(x2); auto y_vec2 = Buffer<T>(y2); auto a_mat2 = Buffer<T>(a2); auto b_mat2 = Buffer<T>(b2); auto c_mat2 = Buffer<T>(c2); + auto dot2 = Buffer<T>(d2); // Runs the two routines - auto status1 = run_reference_(args, Buffers<T>{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1}, queue_); - auto status2 = run_routine_(args, Buffers<T>{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2}, queue_); + auto buffers1 = Buffers<T>{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, dot1}; + auto buffers2 = Buffers<T>{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, dot2}; + auto status1 = run_reference_(args, buffers1, queue_); + auto status2 = run_routine_(args, buffers2, queue_); // Tests for equality of the two status codes TestErrorCodes(status1, status2, args); diff --git a/test/correctness/testblas.h b/test/correctness/testblas.h index 8a86c65e..9e1d110c 100644 --- a/test/correctness/testblas.h +++ b/test/correctness/testblas.h @@ -90,6 +90,7 @@ class TestBlas: public Tester<T,U> { std::vector<T> a_source_; std::vector<T> b_source_; std::vector<T> c_source_; + std::vector<T> dot_source_; // The routine-specific functions passed to the tester Routine run_routine_; @@ -136,6 +137,7 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name auto a_offsets = std::vector<size_t>{args.a_offset}; auto b_offsets = std::vector<size_t>{args.b_offset}; auto c_offsets = std::vector<size_t>{args.c_offset}; + auto dot_offsets = std::vector<size_t>{args.dot_offset}; auto alphas = std::vector<U>{args.alpha}; auto betas = std::vector<U>{args.beta}; auto x_sizes = std::vector<size_t>{args.x_size}; @@ -170,6 +172,7 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name if (option == kArgAOffset) { a_offsets = tester.kOffsets; } if (option == kArgBOffset) { b_offsets = tester.kOffsets; } if (option == kArgCOffset) { c_offsets = tester.kOffsets; } + if (option == kArgDotOffset) { dot_offsets = tester.kOffsets; } if (option == kArgAlpha) { alphas = tester.kAlphaValues; } if (option == kArgBeta) { betas = tester.kBetaValues; } @@ -204,10 +207,12 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name for (auto &b_offset: b_offsets) { r_args.b_offset = b_offset; for (auto &c_ld: c_lds) { r_args.c_ld = c_ld; for (auto &c_offset: c_offsets) { r_args.c_offset = c_offset; - for (auto &alpha: alphas) { r_args.alpha = alpha; - for (auto &beta: betas) { r_args.beta = beta; - C::SetSizes(r_args); - regular_test_vector.push_back(r_args); + for (auto &dot_offset: dot_offsets) { r_args.dot_offset = dot_offset; + for (auto &alpha: alphas) { r_args.alpha = alpha; + for (auto &beta: betas) { r_args.beta = beta; + C::SetSizes(r_args); + regular_test_vector.push_back(r_args); + } } } } diff --git a/test/correctness/tester.cc b/test/correctness/tester.cc index a52142c4..f792925e 100644 --- a/test/correctness/tester.cc +++ b/test/correctness/tester.cc @@ -148,6 +148,7 @@ void Tester<T,U>::TestEnd() { if (o == kArgAOffset) { fprintf(stdout, "%s=%lu ", kArgAOffset, entry.args.a_offset);} if (o == kArgBOffset) { fprintf(stdout, "%s=%lu ", kArgBOffset, entry.args.b_offset);} if (o == kArgCOffset) { fprintf(stdout, "%s=%lu ", kArgCOffset, entry.args.c_offset);} + if (o == kArgDotOffset){ fprintf(stdout, "%s=%lu ", kArgDotOffset, entry.args.dot_offset);} } fprintf(stdout, "\n"); } diff --git a/test/performance/client.cc b/test/performance/client.cc index 893bb55d..9faa4dca 100644 --- a/test/performance/client.cc +++ b/test/performance/client.cc @@ -42,7 +42,7 @@ template <typename T, typename U> Arguments<U> Client<T,U>::ParseArguments(int argc, char *argv[], const GetMetric default_a_ld, const GetMetric default_b_ld, const GetMetric default_c_ld) { auto args = Arguments<U>{}; - auto help = std::string{"Options given/available:\n"}; + auto help = std::string{"\n* Options given/available:\n"}; // These are the options which are not for every client: they are optional for (auto &o: options_) { @@ -74,6 +74,9 @@ Arguments<U> Client<T,U>::ParseArguments(int argc, char *argv[], const GetMetric if (o == kArgBOffset) { args.b_offset = GetArgument(argc, argv, help, kArgBOffset, size_t{0}); } if (o == kArgCOffset) { args.c_offset = GetArgument(argc, argv, help, kArgCOffset, size_t{0}); } + // Dot arguments + if (o == kArgDotOffset) { args.dot_offset = GetArgument(argc, argv, help, kArgDotOffset, size_t{0}); } + // Scalar values if (o == kArgAlpha) { args.alpha = GetArgument(argc, argv, help, kArgAlpha, GetScalar<U>()); } if (o == kArgBeta) { args.beta = GetArgument(argc, argv, help, kArgBeta, GetScalar<U>()); } @@ -128,11 +131,13 @@ void Client<T,U>::PerformanceTest(Arguments<U> &args, const SetMetric set_sizes) std::vector<T> a_source(args.a_size); std::vector<T> b_source(args.b_size); std::vector<T> c_source(args.c_size); + std::vector<T> dot_source(args.dot_size); PopulateVector(x_source); PopulateVector(y_source); PopulateVector(a_source); PopulateVector(b_source); PopulateVector(c_source); + PopulateVector(dot_source); // Creates the matrices on the device auto x_vec = Buffer<T>(context, args.x_size); @@ -140,12 +145,14 @@ void Client<T,U>::PerformanceTest(Arguments<U> &args, const SetMetric set_sizes) auto a_mat = Buffer<T>(context, args.a_size); auto b_mat = Buffer<T>(context, args.b_size); auto c_mat = Buffer<T>(context, args.c_size); + auto dot = Buffer<T>(context, args.dot_size); x_vec.Write(queue, args.x_size, x_source); y_vec.Write(queue, args.y_size, y_source); a_mat.Write(queue, args.a_size, a_source); b_mat.Write(queue, args.b_size, b_source); c_mat.Write(queue, args.c_size, c_source); - auto buffers = Buffers<T>{x_vec, y_vec, a_mat, b_mat, c_mat}; + dot.Write(queue, args.dot_size, dot_source); + auto buffers = Buffers<T>{x_vec, y_vec, a_mat, b_mat, c_mat, dot}; // Runs the routines and collects the timings auto ms_clblast = TimedExecution(args.num_runs, args, buffers, queue, run_routine_, "CLBlast"); @@ -236,6 +243,7 @@ void Client<T,U>::PrintTableRow(const Arguments<U>& args, const double ms_clblas else if (o == kArgAOffset) { integers.push_back(args.a_offset); } else if (o == kArgBOffset) { integers.push_back(args.b_offset); } else if (o == kArgCOffset) { integers.push_back(args.c_offset); } + else if (o == kArgDotOffset) {integers.push_back(args.dot_offset); } } auto strings = std::vector<std::string>{}; for (auto &o: options_) { 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/xdot.cc b/test/performance/routines/level1/xdot.cc new file mode 100644 index 00000000..c82547da --- /dev/null +++ b/test/performance/routines/level1/xdot.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 Xdot command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/level1/xdot.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::TestXdot<float>, float, float>(argc, argv); break; + case clblast::Precision::kDouble: + clblast::RunClient<clblast::TestXdot<double>, double, double>(argc, argv); break; + case clblast::Precision::kComplexSingle: + throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kComplexDouble: + throw std::runtime_error("Unsupported precision mode"); + } + return 0; +} + +// ================================================================================================= diff --git a/test/performance/routines/level1/xdotc.cc b/test/performance/routines/level1/xdotc.cc new file mode 100644 index 00000000..327975d8 --- /dev/null +++ b/test/performance/routines/level1/xdotc.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 Xdotc command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/level1/xdotc.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: + throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kDouble: + throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kComplexSingle: + clblast::RunClient<clblast::TestXdotc<float2>, float2, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXdotc<double2>, double2, double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/test/performance/routines/level1/xdotu.cc b/test/performance/routines/level1/xdotu.cc new file mode 100644 index 00000000..622ffb8e --- /dev/null +++ b/test/performance/routines/level1/xdotu.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 Xdotu command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/level1/xdotu.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: + throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kDouble: + throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kComplexSingle: + clblast::RunClient<clblast::TestXdotu<float2>, float2, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXdotu<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/xdot.h b/test/routines/level1/xdot.h new file mode 100644 index 00000000..bfcfdaff --- /dev/null +++ b/test/routines/level1/xdot.h @@ -0,0 +1,123 @@ + +// ================================================================================================= +// 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 Xdot 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_XDOT_H_ +#define CLBLAST_TEST_ROUTINES_XDOT_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 TestXdot { + 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, kArgDotOffset}; + } + + // 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; + } + static size_t GetSizeDot(const Arguments<T> &args) { + return 1 + args.dot_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); + args.dot_size = GetSizeDot(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 = Dot<T>(args.n, + buffers.dot(), args.dot_offset, + 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 = clblasXdot<T>(args.n, + buffers.dot(), args.dot_offset, + buffers.x_vec(), args.x_offset, args.x_inc, + buffers.y_vec(), args.y_offset, args.y_inc, + 1, &queue_plain, 0, nullptr, &event); + 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.dot_size, static_cast<T>(0)); + buffers.dot.Read(queue, args.dot_size, result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t ResultID2(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t GetResultIndex(const Arguments<T> &args, const size_t, const size_t) { + return args.dot_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<T> &args) { + return 2 * args.n; + } + static size_t GetBytes(const Arguments<T> &args) { + return ((2 * args.n) + 1) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XDOT_H_ +#endif diff --git a/test/routines/level1/xdotc.h b/test/routines/level1/xdotc.h new file mode 100644 index 00000000..e403ba4c --- /dev/null +++ b/test/routines/level1/xdotc.h @@ -0,0 +1,123 @@ + +// ================================================================================================= +// 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 Xdotc 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_XDOTC_H_ +#define CLBLAST_TEST_ROUTINES_XDOTC_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 TestXdotc { + 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, kArgDotOffset}; + } + + // 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; + } + static size_t GetSizeDot(const Arguments<T> &args) { + return 1 + args.dot_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); + args.dot_size = GetSizeDot(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 = Dotc<T>(args.n, + buffers.dot(), args.dot_offset, + 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 = clblasXdotc<T>(args.n, + buffers.dot(), args.dot_offset, + buffers.x_vec(), args.x_offset, args.x_inc, + buffers.y_vec(), args.y_offset, args.y_inc, + 1, &queue_plain, 0, nullptr, &event); + 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.dot_size, static_cast<T>(0)); + buffers.dot.Read(queue, args.dot_size, result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t ResultID2(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t GetResultIndex(const Arguments<T> &args, const size_t, const size_t) { + return args.dot_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<T> &args) { + return 2 * args.n; + } + static size_t GetBytes(const Arguments<T> &args) { + return ((2 * args.n) + 1) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XDOTC_H_ +#endif diff --git a/test/routines/level1/xdotu.h b/test/routines/level1/xdotu.h new file mode 100644 index 00000000..8b2c65a8 --- /dev/null +++ b/test/routines/level1/xdotu.h @@ -0,0 +1,123 @@ + +// ================================================================================================= +// 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 Xdotu 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_XDOTU_H_ +#define CLBLAST_TEST_ROUTINES_XDOTU_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 TestXdotu { + 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, kArgDotOffset}; + } + + // 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; + } + static size_t GetSizeDot(const Arguments<T> &args) { + return 1 + args.dot_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); + args.dot_size = GetSizeDot(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 = Dotu<T>(args.n, + buffers.dot(), args.dot_offset, + 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 = clblasXdotu<T>(args.n, + buffers.dot(), args.dot_offset, + buffers.x_vec(), args.x_offset, args.x_inc, + buffers.y_vec(), args.y_offset, args.y_inc, + 1, &queue_plain, 0, nullptr, &event); + 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.dot_size, static_cast<T>(0)); + buffers.dot.Read(queue, args.dot_size, result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t ResultID2(const Arguments<T> &) { return 1; } // N/A for this routine + static size_t GetResultIndex(const Arguments<T> &args, const size_t, const size_t) { + return args.dot_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<T> &args) { + return 2 * args.n; + } + static size_t GetBytes(const Arguments<T> &args) { + return ((2 * args.n) + 1) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XDOTU_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..fcf1a918 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, @@ -71,6 +215,135 @@ clblasStatus clblasXaxpy( num_queues, queues, num_wait_events, wait_events, events); } +// Forwards the clBLAS calls for SDOT/DDOT +template <typename T> clblasStatus clblasXdot( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events); +template <> clblasStatus clblasXdot<float>( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto queue = Queue(queues[0]); + auto context = queue.GetContext(); + auto scratch_buffer = Buffer<float>(context, n); + return clblasSdot(n, + dot_buffer, dot_offset, + x_buffer, x_offset, static_cast<int>(x_inc), + y_buffer, y_offset, static_cast<int>(y_inc), + scratch_buffer(), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXdot<double>( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto queue = Queue(queues[0]); + auto context = queue.GetContext(); + auto scratch_buffer = Buffer<double>(context, n); + return clblasDdot(n, + dot_buffer, dot_offset, + x_buffer, x_offset, static_cast<int>(x_inc), + y_buffer, y_offset, static_cast<int>(y_inc), + scratch_buffer(), + num_queues, queues, num_wait_events, wait_events, events); +} + +// Forwards the clBLAS calls for CDOTU/ZDOTU +template <typename T> clblasStatus clblasXdotu( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events); +template <> clblasStatus clblasXdotu<float2>( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto queue = Queue(queues[0]); + auto context = queue.GetContext(); + auto scratch_buffer = Buffer<float2>(context, n); + return clblasCdotu(n, + dot_buffer, dot_offset, + x_buffer, x_offset, static_cast<int>(x_inc), + y_buffer, y_offset, static_cast<int>(y_inc), + scratch_buffer(), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXdotu<double2>( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto queue = Queue(queues[0]); + auto context = queue.GetContext(); + auto scratch_buffer = Buffer<double2>(context, n); + return clblasZdotu(n, + dot_buffer, dot_offset, + x_buffer, x_offset, static_cast<int>(x_inc), + y_buffer, y_offset, static_cast<int>(y_inc), + scratch_buffer(), + num_queues, queues, num_wait_events, wait_events, events); +} + +// Forwards the clBLAS calls for CDOTC/ZDOTC +template <typename T> clblasStatus clblasXdotc( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events); +template <> clblasStatus clblasXdotc<float2>( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto queue = Queue(queues[0]); + auto context = queue.GetContext(); + auto scratch_buffer = Buffer<float2>(context, n); + return clblasCdotc(n, + dot_buffer, dot_offset, + x_buffer, x_offset, static_cast<int>(x_inc), + y_buffer, y_offset, static_cast<int>(y_inc), + scratch_buffer(), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> clblasStatus clblasXdotc<double2>( + const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto queue = Queue(queues[0]); + auto context = queue.GetContext(); + auto scratch_buffer = Buffer<double2>(context, n); + return clblasZdotc(n, + dot_buffer, dot_offset, + x_buffer, x_offset, static_cast<int>(x_inc), + y_buffer, y_offset, static_cast<int>(y_inc), + scratch_buffer(), + num_queues, queues, num_wait_events, wait_events, events); +} + // ================================================================================================= // BLAS level-2 (matrix-vector) routines |