diff options
57 files changed, 3143 insertions, 360 deletions
@@ -2,8 +2,12 @@ Development version (next release) - Re-organized test/client infrastructure to avoid code duplication - Added level-3 routines: + * CHEMM/ZHEMM * SSYRK/DSYRK/CSYRK/ZSYRK + * CHERK/ZHERK * SSYR2K/DSYR2K/CSYR2K/ZSYR2K + * CHER2K/ZHER2K + * STRMM/DTRMM/CTRMM/ZTRMM Version 0.2.0 - Added support for complex conjugate transpose diff --git a/CMakeLists.txt b/CMakeLists.txt index 60b1aaed..1b2c5657 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -98,7 +98,7 @@ set(SAMPLE_PROGRAMS sgemm) set(ROUTINES xaxpy xgemv - xgemm xsymm xsyrk xsyr2k) + xgemm xsymm xhemm xsyrk xherk xsyr2k xher2k xtrmm) # ================================================================================================== @@ -177,12 +177,12 @@ CLBlast is in active development and currently does not support the full set of | ---------|---|---|---|---|---------| | xGEMM | ✔ | ✔ | ✔ | ✔ | | | xSYMM | ✔ | ✔ | ✔ | ✔ | | -| xHEMM | - | - | | | | +| xHEMM | - | - | ✔ | ✔ | | | xSYRK | ✔ | ✔ | ✔ | ✔ | | -| xHERK | - | - | | | | +| xHERK | - | - | ✔ | ✔ | | | xSYR2K | ✔ | ✔ | ✔ | ✔ | | -| xHER2K | - | - | | | | -| xTRMM | | | | | | +| xHER2K | - | - | ✔ | ✔ | | +| xTRMM | ✔ | ✔ | ✔ | ✔ | | | xTRSM | | | | | | diff --git a/include/clblast.h b/include/clblast.h index da504a0b..80ea1707 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -75,6 +75,7 @@ enum class Layout { kRowMajor, kColMajor }; enum class Transpose { kNo, kYes, kConjugate }; enum class Side { kLeft, kRight }; enum class Triangle { kUpper, kLower }; +enum class Diagonal { kUnit, kNonUnit }; // Precision scoped enum (values in bits) enum class Precision { kHalf = 16, kSingle = 32, kDouble = 64, @@ -95,7 +96,7 @@ StatusCode Axpy(const size_t n, const T alpha, // Templated-precision generalized matrix-vector multiplication: SGEMV/DGEMV/CGEMV/ZGEMV template <typename T> -StatusCode Gemv(const Layout layout, const Transpose transpose_a, +StatusCode Gemv(const Layout layout, const Transpose a_transpose, const size_t m, const size_t n, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, @@ -109,7 +110,7 @@ StatusCode Gemv(const Layout layout, const Transpose transpose_a, // Templated-precision generalized matrix-matrix multiplication: SGEMM/DGEMM/CGEMM/ZGEMM template <typename T> -StatusCode Gemm(const Layout layout, const Transpose transpose_a, const Transpose transpose_b, +StatusCode Gemm(const Layout layout, const Transpose a_transpose, const Transpose b_transpose, const size_t m, const size_t n, const size_t k, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, @@ -129,9 +130,30 @@ StatusCode Symm(const Layout layout, const Side side, const Triangle triangle, cl_mem c_buffer, const size_t c_offset, const size_t c_ld, cl_command_queue* queue, cl_event* event); +// Templated-precision hermitian matrix-matrix multiplication: CHEMM/ZHEMM +template <typename T> +StatusCode Hemm(const Layout layout, const Side side, const Triangle triangle, + const size_t m, const size_t n, + const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, + const T beta, + cl_mem c_buffer, const size_t c_offset, const size_t c_ld, + cl_command_queue* queue, cl_event* event); + // Templated-precision rank-K update of a symmetric matrix: SSYRK/DSYRK/CSYRK/ZSYRK template <typename T> -StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose transpose_a, +StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose a_transpose, + const size_t n, const size_t k, + const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + const T beta, + cl_mem c_buffer, const size_t c_offset, const size_t c_ld, + cl_command_queue* queue, cl_event* event); + +// Templated-precision rank-K update of a hermitian matrix: CHERK/ZHERK +template <typename T> +StatusCode Herk(const Layout layout, const Triangle triangle, const Transpose a_transpose, const size_t n, const size_t k, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, @@ -141,7 +163,7 @@ StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose tr // Templated-precision rank-2K update of a symmetric matrix: SSYR2K/DSYR2K/CSYR2K/ZSYR2K template <typename T> -StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose transpose_ab, +StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose ab_transpose, const size_t n, const size_t k, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, @@ -150,6 +172,39 @@ StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose t cl_mem c_buffer, const size_t c_offset, const size_t c_ld, cl_command_queue* queue, cl_event* event); +// Templated-precision rank-2K update of a hermitian matrix: CHER2K/ZHER2K +template <typename T, typename U> +StatusCode Her2k(const Layout layout, const Triangle triangle, const Transpose ab_transpose, + const size_t n, const size_t k, + const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, + const U beta, + cl_mem c_buffer, const size_t c_offset, const size_t c_ld, + cl_command_queue* queue, cl_event* event); + +// Templated-precision triangular matrix-matrix multiplication: STRMM/DTRMM/CTRMM/ZTRMM +template <typename T> +StatusCode Trmm(const Layout layout, const Side side, const Triangle triangle, + const Transpose a_transpose, const Diagonal diagonal, + const size_t m, const size_t n, + const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + cl_mem b_buffer, const size_t b_offset, const size_t b_ld, + cl_command_queue* queue, cl_event* event); + +// Templated-precision matrix equation solver: STRSM/DTRSM/CTRSM/ZTRSM +/* +template <typename T> +StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, + const Transpose a_transpose, const Diagonal diagonal, + const size_t m, const size_t n, + const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + cl_mem b_buffer, const size_t b_offset, const size_t b_ld, + cl_command_queue* queue, cl_event* event); +*/ + // ================================================================================================= } // namespace clblast diff --git a/include/internal/routine.h b/include/internal/routine.h index d2cee52c..acc9a9c8 100644 --- a/include/internal/routine.h +++ b/include/internal/routine.h @@ -93,6 +93,7 @@ class Routine { const Buffer &dest, const bool do_transpose, const bool do_conjugate, const bool pad, const bool upper, const bool lower, + const bool diagonal_imag_zero, const Program &program); // Queries the cache and retrieve either a matching program or a boolean whether a match exists. diff --git a/include/internal/routines/xhemm.h b/include/internal/routines/xhemm.h new file mode 100644 index 00000000..1b1a0dfa --- /dev/null +++ b/include/internal/routines/xhemm.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 Xhemm routine. It is based on the generalized matrix multiplication +// routine (Xgemm). The implementation is very similar to the Xsymm routine. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XHEMM_H_ +#define CLBLAST_ROUTINES_XHEMM_H_ + +#include "internal/routines/xgemm.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xhemm: public Xgemm<T> { + public: + + // Uses several variables from the Routine class + using Routine::db_; + using Routine::context_; + + // Uses several helper functions from the Routine class + using Routine::RunKernel; + using Routine::ErrorIn; + using Routine::TestMatrixA; + using Routine::GetProgramFromCache; + + // Uses the regular Xgemm routine + using Xgemm<T>::DoGemm; + + // Constructor + Xhemm(CommandQueue &queue, Event &event); + + // Templated-precision implementation of the routine + StatusCode DoHemm(const Layout layout, const Side side, const Triangle triangle, + const size_t m, const size_t n, + const T alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const Buffer &b_buffer, const size_t b_offset, const size_t b_ld, + const T beta, + const Buffer &c_buffer, const size_t c_offset, const size_t c_ld); +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XHEMM_H_ +#endif diff --git a/include/internal/routines/xher2k.h b/include/internal/routines/xher2k.h new file mode 100644 index 00000000..1836a812 --- /dev/null +++ b/include/internal/routines/xher2k.h @@ -0,0 +1,48 @@ + +// ================================================================================================= +// 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 Xher2k routine. The precision is implemented using the template argument +// 'T', whereas the alpha/beta arguments are of type 'U'. The implementation is very similar to the +// Xsyr2k routine. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XHER2K_H_ +#define CLBLAST_ROUTINES_XHER2K_H_ + +#include "internal/routine.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T, typename U> +class Xher2k: public Routine { + public: + Xher2k(CommandQueue &queue, Event &event); + + // Templated-precision implementation of the routine + StatusCode DoHer2k(const Layout layout, const Triangle triangle, const Transpose ab_transpose, + const size_t n, const size_t k, + const T alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const Buffer &b_buffer, const size_t b_offset, const size_t b_ld, + const U beta, + const Buffer &c_buffer, const size_t c_offset, const size_t c_ld); + + private: + // Static variable to get the precision + const static Precision precision_; +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XHER2K_H_ +#endif diff --git a/include/internal/routines/xherk.h b/include/internal/routines/xherk.h new file mode 100644 index 00000000..9b361254 --- /dev/null +++ b/include/internal/routines/xherk.h @@ -0,0 +1,47 @@ + +// ================================================================================================= +// 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 Xherk routine. The precision is implemented using the template argument +// 'T', whereas the alpha/beta arguments are of type 'U'. The implementation is very similar to the +// Xsyrk routine. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XHERK_H_ +#define CLBLAST_ROUTINES_XHERK_H_ + +#include "internal/routine.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T, typename U> +class Xherk: public Routine { + public: + Xherk(CommandQueue &queue, Event &event); + + // Templated-precision implementation of the routine + StatusCode DoHerk(const Layout layout, const Triangle triangle, const Transpose a_transpose, + const size_t n, const size_t k, + const U alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const U beta, + const Buffer &c_buffer, const size_t c_offset, const size_t c_ld); + + private: + // Static variable to get the precision + const static Precision precision_; +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XHERK_H_ +#endif diff --git a/include/internal/routines/xtrmm.h b/include/internal/routines/xtrmm.h new file mode 100644 index 00000000..af9f0266 --- /dev/null +++ b/include/internal/routines/xtrmm.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 Xtrmm routine. The implementation is based on first transforming the +// upper/lower unit/non-unit triangular matrix into a regular matrix and then calling the GEMM +// routine. Therefore, this class inherits from the Xgemm class. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XTRMM_H_ +#define CLBLAST_ROUTINES_XTRMM_H_ + +#include "internal/routines/xgemm.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template <typename T> +class Xtrmm: public Xgemm<T> { + public: + + // Uses several variables from the Routine class + using Routine::db_; + using Routine::context_; + + // Uses several helper functions from the Routine class + using Routine::RunKernel; + using Routine::ErrorIn; + using Routine::TestMatrixA; + using Routine::GetProgramFromCache; + + // Uses the regular Xgemm routine + using Xgemm<T>::DoGemm; + + // Constructor + Xtrmm(CommandQueue &queue, Event &event); + + // Templated-precision implementation of the routine + StatusCode DoTrmm(const Layout layout, const Side side, const Triangle triangle, + const Transpose a_transpose, const Diagonal diagonal, + const size_t m, const size_t n, + const T alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const Buffer &b_buffer, const size_t b_offset, const size_t b_ld); +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XTRMM_H_ +#endif diff --git a/include/internal/utilities.h b/include/internal/utilities.h index 93cd509e..60d70eae 100644 --- a/include/internal/utilities.h +++ b/include/internal/utilities.h @@ -46,6 +46,7 @@ constexpr auto kArgATransp = "transA"; constexpr auto kArgBTransp = "transB"; constexpr auto kArgSide = "side"; constexpr auto kArgTriangle = "triangle"; +constexpr auto kArgDiagonal = "diagonal"; constexpr auto kArgXInc = "incx"; constexpr auto kArgYInc = "incy"; constexpr auto kArgXOffset = "offx"; @@ -93,6 +94,7 @@ struct Arguments { Transpose b_transpose = Transpose::kNo; Side side = Side::kLeft; Triangle triangle = Triangle::kUpper; + Diagonal diagonal = Diagonal::kUnit; size_t x_inc = 1; size_t y_inc = 1; size_t x_offset = 0; diff --git a/src/clblast.cc b/src/clblast.cc index b8aa1e39..23046b01 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -26,8 +26,12 @@ // BLAS level-3 includes #include "internal/routines/xgemm.h" #include "internal/routines/xsymm.h" +#include "internal/routines/xhemm.h" #include "internal/routines/xsyrk.h" +#include "internal/routines/xherk.h" #include "internal/routines/xsyr2k.h" +#include "internal/routines/xher2k.h" +#include "internal/routines/xtrmm.h" namespace clblast { // ================================================================================================= @@ -76,7 +80,7 @@ template StatusCode Axpy<double2>(const size_t, const double2, // GEMV template <typename T> -StatusCode Gemv(const Layout layout, const Transpose transpose_a, +StatusCode Gemv(const Layout layout, const Transpose a_transpose, const size_t m, const size_t n, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, const T beta, @@ -94,7 +98,7 @@ StatusCode Gemv(const Layout layout, const Transpose transpose_a, if (status != StatusCode::kSuccess) { return status; } // Runs the routine - return routine.DoGemv(layout, transpose_a, m, n, alpha, + return routine.DoGemv(layout, a_transpose, m, n, alpha, Buffer(a_buffer), a_offset, a_ld, Buffer(x_buffer), x_offset, x_inc, beta, Buffer(y_buffer), y_offset, y_inc); @@ -129,7 +133,7 @@ template StatusCode Gemv<double2>(const Layout, const Transpose, // GEMM template <typename T> -StatusCode Gemm(const Layout layout, const Transpose transpose_a, const Transpose transpose_b, +StatusCode Gemm(const Layout layout, const Transpose a_transpose, const Transpose b_transpose, const size_t m, const size_t n, const size_t k, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, const T beta, @@ -155,7 +159,7 @@ StatusCode Gemm(const Layout layout, const Transpose transpose_a, const Transpos if (status != StatusCode::kSuccess) { return status; } // Runs the routine - return routine.DoGemm(layout, transpose_a, transpose_b, m, n, k, alpha, + return routine.DoGemm(layout, a_transpose, b_transpose, m, n, k, alpha, Buffer(a_buffer), a_offset, a_ld, Buffer(b_buffer), b_offset, b_ld, beta, Buffer(c_buffer), c_offset, c_ld); @@ -247,9 +251,57 @@ template StatusCode Symm<double2>(const Layout, const Side, const Triangle, // ================================================================================================= +// HEMM +template <typename T> +StatusCode Hemm(const Layout layout, const Side side, const Triangle triangle, + const size_t m, const size_t n, const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, const T beta, + cl_mem c_buffer, const size_t c_offset, const size_t c_ld, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = CommandQueue(*queue); + auto event_cpp = Event(*event); + auto routine = Xhemm<T>(queue_cpp, event_cpp); + + // Loads the kernel source-code as an include (C++11 raw string literal) + std::string common_source1 = + #include "kernels/copy.opencl" + std::string common_source2 = + #include "kernels/pad.opencl" + std::string common_source3 = + #include "kernels/transpose.opencl" + std::string common_source4 = + #include "kernels/padtranspose.opencl" + std::string kernel_source = + #include "kernels/xgemm.opencl" + auto status = routine.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + + kernel_source); + if (status != StatusCode::kSuccess) { return status; } + + // Runs the routine + return routine.DoHemm(layout, side, triangle, m, n, alpha, + Buffer(a_buffer), a_offset, a_ld, + Buffer(b_buffer), b_offset, b_ld, beta, + Buffer(c_buffer), c_offset, c_ld); +} +template StatusCode Hemm<float2>(const Layout, const Side, const Triangle, + const size_t, const size_t, const float2, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, const float2, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); +template StatusCode Hemm<double2>(const Layout, const Side, const Triangle, + const size_t, const size_t, const double2, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, const double2, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); + +// ================================================================================================= + // SYRK template <typename T> -StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose transpose_a, +StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose a_transpose, const size_t n, const size_t k, const T alpha, const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, const T beta, cl_mem c_buffer, const size_t c_offset, const size_t c_ld, @@ -274,7 +326,7 @@ StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose tr if (status != StatusCode::kSuccess) { return status; } // Runs the routine - return routine.DoSyrk(layout, triangle, transpose_a, n, k, alpha, + return routine.DoSyrk(layout, triangle, a_transpose, n, k, alpha, Buffer(a_buffer), a_offset, a_ld, beta, Buffer(c_buffer), c_offset, c_ld); } @@ -301,16 +353,60 @@ template StatusCode Syrk<double2>(const Layout, const Triangle, const Transpose, // ================================================================================================= -// SYR2K +// HERK template <typename T> -StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose transpose_ab, +StatusCode Herk(const Layout layout, const Triangle triangle, const Transpose a_transpose, const size_t n, const size_t k, const T alpha, - const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, - const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, const T beta, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, const T beta, cl_mem c_buffer, const size_t c_offset, const size_t c_ld, cl_command_queue* queue, cl_event* event) { auto queue_cpp = CommandQueue(*queue); auto event_cpp = Event(*event); + auto routine = Xherk<std::complex<T>,T>(queue_cpp, event_cpp); + + // Loads the kernel source-code as an include (C++11 raw string literal) + std::string common_source1 = + #include "kernels/copy.opencl" + std::string common_source2 = + #include "kernels/pad.opencl" + std::string common_source3 = + #include "kernels/transpose.opencl" + std::string common_source4 = + #include "kernels/padtranspose.opencl" + std::string kernel_source = + #include "kernels/xgemm.opencl" + auto status = routine.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + + kernel_source); + if (status != StatusCode::kSuccess) { return status; } + + // Runs the routine + return routine.DoHerk(layout, triangle, a_transpose, n, k, alpha, + Buffer(a_buffer), a_offset, a_ld, beta, + Buffer(c_buffer), c_offset, c_ld); +} +template StatusCode Herk<float>(const Layout, const Triangle, const Transpose, + const size_t, const size_t, const float, + const cl_mem, const size_t, const size_t, const float, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); +template StatusCode Herk<double>(const Layout, const Triangle, const Transpose, + const size_t, const size_t, const double, + const cl_mem, const size_t, const size_t, const double, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); + +// ================================================================================================= + +// SYR2K +template <typename T> +StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose ab_transpose, + const size_t n, const size_t k, const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, const T beta, + cl_mem c_buffer, const size_t c_offset, const size_t c_ld, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = CommandQueue(*queue); + auto event_cpp = Event(*event); auto routine = Xsyr2k<T>(queue_cpp, event_cpp); // Loads the kernel source-code as an include (C++11 raw string literal) @@ -329,7 +425,7 @@ StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose t if (status != StatusCode::kSuccess) { return status; } // Runs the routine - return routine.DoSyr2k(layout, triangle, transpose_ab, n, k, alpha, + return routine.DoSyr2k(layout, triangle, ab_transpose, n, k, alpha, Buffer(a_buffer), a_offset, a_ld, Buffer(b_buffer), b_offset, b_ld, beta, Buffer(c_buffer), c_offset, c_ld); @@ -360,4 +456,173 @@ template StatusCode Syr2k<double2>(const Layout, const Triangle, const Transpose cl_command_queue*, cl_event*); // ================================================================================================= + +// SYR2K +template <typename T, typename U> +StatusCode Her2k(const Layout layout, const Triangle triangle, const Transpose ab_transpose, + const size_t n, const size_t k, const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, const U beta, + cl_mem c_buffer, const size_t c_offset, const size_t c_ld, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = CommandQueue(*queue); + auto event_cpp = Event(*event); + auto routine = Xher2k<T,U>(queue_cpp, event_cpp); + + // Loads the kernel source-code as an include (C++11 raw string literal) + std::string common_source1 = + #include "kernels/copy.opencl" + std::string common_source2 = + #include "kernels/pad.opencl" + std::string common_source3 = + #include "kernels/transpose.opencl" + std::string common_source4 = + #include "kernels/padtranspose.opencl" + std::string kernel_source = + #include "kernels/xgemm.opencl" + auto status = routine.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + + kernel_source); + if (status != StatusCode::kSuccess) { return status; } + + // Runs the routine + return routine.DoHer2k(layout, triangle, ab_transpose, n, k, alpha, + Buffer(a_buffer), a_offset, a_ld, + Buffer(b_buffer), b_offset, b_ld, beta, + Buffer(c_buffer), c_offset, c_ld); +} +template StatusCode Her2k<float2,float>(const Layout, const Triangle, const Transpose, + const size_t, const size_t, const float2, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, const float, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); +template StatusCode Her2k<double2,double>(const Layout, const Triangle, const Transpose, + const size_t, const size_t, const double2, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, const double, + cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); + +// ================================================================================================= + +// TRMM +template <typename T> +StatusCode Trmm(const Layout layout, const Side side, const Triangle triangle, + const Transpose a_transpose, const Diagonal diagonal, + const size_t m, const size_t n, + const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + cl_mem b_buffer, const size_t b_offset, const size_t b_ld, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = CommandQueue(*queue); + auto event_cpp = Event(*event); + auto routine = Xtrmm<T>(queue_cpp, event_cpp); + + // Loads the kernel source-code as an include (C++11 raw string literal) + std::string common_source1 = + #include "kernels/copy.opencl" + std::string common_source2 = + #include "kernels/pad.opencl" + std::string common_source3 = + #include "kernels/transpose.opencl" + std::string common_source4 = + #include "kernels/padtranspose.opencl" + std::string kernel_source = + #include "kernels/xgemm.opencl" + auto status = routine.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + + kernel_source); + if (status != StatusCode::kSuccess) { return status; } + + // Runs the routine + return routine.DoTrmm(layout, side, triangle, a_transpose, diagonal, m, n, alpha, + Buffer(a_buffer), a_offset, a_ld, + Buffer(b_buffer), b_offset, b_ld); +} +template StatusCode Trmm<float>(const Layout, const Side, const Triangle, + const Transpose, const Diagonal, + const size_t, 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 Trmm<double>(const Layout, const Side, const Triangle, + const Transpose, const Diagonal, + const size_t, 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 Trmm<float2>(const Layout, const Side, const Triangle, + const Transpose, const Diagonal, + const size_t, 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 Trmm<double2>(const Layout, const Side, const Triangle, + const Transpose, const Diagonal, + const size_t, 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*); + +// ================================================================================================= + +// TRSM +/* +template <typename T> +StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, + const Transpose a_transpose, const Diagonal diagonal, + const size_t m, const size_t n, + const T alpha, + const cl_mem a_buffer, const size_t a_offset, const size_t a_ld, + cl_mem b_buffer, const size_t b_offset, const size_t b_ld, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = CommandQueue(*queue); + auto event_cpp = Event(*event); + auto routine = Xtrsm<T>(queue_cpp, event_cpp); + + // Loads the kernel source-code as an include (C++11 raw string literal) + std::string common_source1 = + #include "kernels/copy.opencl" + std::string common_source2 = + #include "kernels/pad.opencl" + std::string common_source3 = + #include "kernels/transpose.opencl" + std::string common_source4 = + #include "kernels/padtranspose.opencl" + std::string kernel_source = + #include "kernels/xgemm.opencl" + auto status = routine.SetUp(common_source1 + common_source2 + common_source3 + common_source4 + + kernel_source); + if (status != StatusCode::kSuccess) { return status; } + + // Runs the routine + return routine.DoTrsm(layout, side, triangle, a_transpose, diagonal, m, n, alpha, + Buffer(a_buffer), a_offset, a_ld, + Buffer(b_buffer), b_offset, b_ld); +} +template StatusCode Trsm<float>(const Layout, const Side, const Triangle, + const Transpose, const Diagonal, + const size_t, 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 Trsm<double>(const Layout, const Side, const Triangle, + const Transpose, const Diagonal, + const size_t, 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 Trsm<float2>(const Layout, const Side, const Triangle, + const Transpose, const Diagonal, + const size_t, 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 Trsm<double2>(const Layout, const Side, const Triangle, + const Transpose, const Diagonal, + const size_t, 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*); +*/ +// ================================================================================================= } // namespace clblast diff --git a/src/kernels/common.opencl b/src/kernels/common.opencl index 818c725f..0d29c7a6 100644 --- a/src/kernels/common.opencl +++ b/src/kernels/common.opencl @@ -39,6 +39,7 @@ R"( typedef float8 real8; typedef float16 real16; #define ZERO 0.0f + #define ONE 1.0f // Double-precision #elif PRECISION == 64 @@ -48,6 +49,7 @@ R"( typedef double8 real8; typedef double16 real16; #define ZERO 0.0 + #define ONE 1.0 // Complex single-precision #elif PRECISION == 3232 @@ -61,6 +63,7 @@ R"( real s8; real s9; real sA; real sB; real sC; real sD; real sE; real sF;} real16; #define ZERO 0.0f + #define ONE 1.0f // Complex Double-precision #elif PRECISION == 6464 @@ -74,6 +77,7 @@ R"( real s8; real s9; real sA; real sB; real sC; real sD; real sE; real sF;} real16; #define ZERO 0.0 + #define ONE 1.0 #endif // ================================================================================================= @@ -88,6 +92,20 @@ R"( #define SetToZero(a) a = ZERO #endif +// Sets a variable to zero (only the imaginary part) +#if PRECISION == 3232 || PRECISION == 6464 + #define ImagToZero(a) a.y = ZERO +#else + #define ImagToZero(a) +#endif + +// Sets a variable to one +#if PRECISION == 3232 || PRECISION == 6464 + #define SetToOne(a) a.x = ONE; a.y = ZERO +#else + #define SetToOne(a) a = ONE +#endif + // Multiply two complex variables (used in the define below) #if PRECISION == 3232 || PRECISION == 6464 #define MulReal(a, b) a.x*b.x - a.y*b.y diff --git a/src/kernels/pad.opencl b/src/kernels/pad.opencl index cce0c746..2791db30 100644 --- a/src/kernels/pad.opencl +++ b/src/kernels/pad.opencl @@ -87,7 +87,8 @@ __kernel void UnPadMatrix(const int src_one, const int src_two, const int dest_one, const int dest_two, const int dest_ld, const int dest_offset, __global real* dest, - const int upper, const int lower) { + const int upper, const int lower, + const int diagonal_imag_zero) { // Loops over the work per thread in both dimensions #pragma unroll @@ -106,7 +107,9 @@ __kernel void UnPadMatrix(const int src_one, const int src_two, // Copies the value into the destination matrix. This is always within bounds of the source // matrix, as we know that the destination matrix is smaller than the source. if (id_two < dest_two && id_one < dest_one) { - dest[id_two*dest_ld + id_one + dest_offset] = src[id_two*src_ld + id_one + src_offset]; + real value = src[id_two*src_ld + id_one + src_offset]; + if (diagonal_imag_zero == 1 && id_one == id_two) { ImagToZero(value); } + dest[id_two*dest_ld + id_one + dest_offset] = value; } } } @@ -135,15 +138,15 @@ __kernel void SymmLowerToSquared(const int src_dim, if (id_two < dest_dim && id_one < dest_dim) { // Loads data from the lower-symmetric matrix - real value; - SetToZero(value); + real result; + SetToZero(result); if (id_two < src_dim && id_one < src_dim) { - if (id_two <= id_one) { value = src[id_two*src_ld + id_one + src_offset]; } - else { value = src[id_one*src_ld + id_two + src_offset]; } + if (id_two <= id_one) { result = src[id_two*src_ld + id_one + src_offset]; } + else { result = src[id_one*src_ld + id_two + src_offset]; } } - // Stores the value in the destination matrix - dest[id_two*dest_ld + id_one + dest_offset] = value; + // Stores the result in the destination matrix + dest[id_two*dest_ld + id_one + dest_offset] = result; } } } @@ -168,15 +171,171 @@ __kernel void SymmUpperToSquared(const int src_dim, if (id_two < dest_dim && id_one < dest_dim) { // Loads data from the upper-symmetric matrix - real value; - SetToZero(value); + real result; + SetToZero(result); if (id_two < src_dim && id_one < src_dim) { - if (id_one <= id_two) { value = src[id_two*src_ld + id_one + src_offset]; } - else { value = src[id_one*src_ld + id_two + src_offset]; } + if (id_one <= id_two) { result = src[id_two*src_ld + id_one + src_offset]; } + else { result = src[id_one*src_ld + id_two + src_offset]; } } - // Stores the value in the destination matrix - dest[id_two*dest_ld + id_one + dest_offset] = value; + // Stores the result in the destination matrix + dest[id_two*dest_ld + id_one + dest_offset] = result; + } + } + } +} + +// ================================================================================================= +#if PRECISION == 3232 || PRECISION == 6464 + +// Kernel to populate a squared hermitian matrix, given that the triangle which holds the data is +// stored as the lower-triangle of the input matrix. This uses the padding kernel's parameters. +__attribute__((reqd_work_group_size(PAD_DIMX, PAD_DIMY, 1))) +__kernel void HermLowerToSquared(const int src_dim, + const int src_ld, const int src_offset, + __global const real* restrict src, + const int dest_dim, + const int dest_ld, const int dest_offset, + __global real* dest) { + + // Loops over the work per thread in both dimensions + #pragma unroll + for (int w_one=0; w_one<PAD_WPTX; ++w_one) { + const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0); + #pragma unroll + for (int w_two=0; w_two<PAD_WPTY; ++w_two) { + const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1); + if (id_two < dest_dim && id_one < dest_dim) { + + // Loads data from the lower-hermitian matrix + real result; + SetToZero(result); + if (id_two < src_dim && id_one < src_dim) { + if (id_two <= id_one) { + result = src[id_two*src_ld + id_one + src_offset]; + if (id_one == id_two) { result.y = ZERO; } + } + else { + result = src[id_one*src_ld + id_two + src_offset]; + COMPLEX_CONJUGATE(result); + } + } + + // Stores the result in the destination matrix + dest[id_two*dest_ld + id_one + dest_offset] = result; + } + } + } +} + +// Same as above, but now the matrix' data is stored in the upper-triangle +__attribute__((reqd_work_group_size(PAD_DIMX, PAD_DIMY, 1))) +__kernel void HermUpperToSquared(const int src_dim, + const int src_ld, const int src_offset, + __global const real* restrict src, + const int dest_dim, + const int dest_ld, const int dest_offset, + __global real* dest) { + + // Loops over the work per thread in both dimensions + #pragma unroll + for (int w_one=0; w_one<PAD_WPTX; ++w_one) { + const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0); + #pragma unroll + for (int w_two=0; w_two<PAD_WPTY; ++w_two) { + const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1); + if (id_two < dest_dim && id_one < dest_dim) { + + // Loads data from the upper-hermitian matrix + real result; + SetToZero(result); + if (id_two < src_dim && id_one < src_dim) { + if (id_one <= id_two) { + result = src[id_two*src_ld + id_one + src_offset]; + if (id_one == id_two) { result.y = ZERO; } + } + else { + result = src[id_one*src_ld + id_two + src_offset]; + COMPLEX_CONJUGATE(result); + } + } + + // Stores the result in the destination matrix + dest[id_two*dest_ld + id_one + dest_offset] = result; + } + } + } +} + +#endif +// ================================================================================================= + +// Kernel to populate a squared triangular matrix, given that the triangle which holds the data is +// stored as the lower-triangle of the input matrix. This uses the padding kernel's parameters. +__attribute__((reqd_work_group_size(PAD_DIMX, PAD_DIMY, 1))) +__kernel void TrmmLowerToSquared(const int src_dim, + const int src_ld, const int src_offset, + __global const real* restrict src, + const int dest_dim, + const int dest_ld, const int dest_offset, + __global real* dest, + const int unit_diagonal) { + + // Loops over the work per thread in both dimensions + #pragma unroll + for (int w_one=0; w_one<PAD_WPTX; ++w_one) { + const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0); + #pragma unroll + for (int w_two=0; w_two<PAD_WPTY; ++w_two) { + const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1); + if (id_two < dest_dim && id_one < dest_dim) { + + // Loads data from the lower-triangular matrix + real result; + SetToZero(result); + if (id_two < src_dim && id_one < src_dim) { + if (id_two <= id_one) { result = src[id_two*src_ld + id_one + src_offset]; } + if (id_two == id_one && unit_diagonal) { SetToOne(result); } + // Else: result is zero + } + + // Stores the result in the destination matrix + dest[id_two*dest_ld + id_one + dest_offset] = result; + } + } + } +} + +// Same as above, but now the matrix' data is stored in the upper-triangle +__attribute__((reqd_work_group_size(PAD_DIMX, PAD_DIMY, 1))) +__kernel void TrmmUpperToSquared(const int src_dim, + const int src_ld, const int src_offset, + __global const real* restrict src, + const int dest_dim, + const int dest_ld, const int dest_offset, + __global real* dest, + const int unit_diagonal) { + + // Loops over the work per thread in both dimensions + #pragma unroll + for (int w_one=0; w_one<PAD_WPTX; ++w_one) { + const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0); + #pragma unroll + for (int w_two=0; w_two<PAD_WPTY; ++w_two) { + const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1); + if (id_two < dest_dim && id_one < dest_dim) { + + // Loads data from the upper-triangular matrix + real result; + SetToZero(result); + if (id_two < src_dim && id_one < src_dim) { + if (id_one <= id_two) { result = src[id_two*src_ld + id_one + src_offset]; } + if (id_one == id_two && unit_diagonal) { SetToOne(result); } + // Else: result is zero + } + + // Stores the result in the destination matrix + dest[id_two*dest_ld + id_one + dest_offset] = result; } } } diff --git a/src/kernels/padtranspose.opencl b/src/kernels/padtranspose.opencl index 7e923392..b2b96aa0 100644 --- a/src/kernels/padtranspose.opencl +++ b/src/kernels/padtranspose.opencl @@ -101,7 +101,8 @@ __kernel void UnPadTransposeMatrix(const int src_one, const int src_two, const int dest_one, const int dest_two, const int dest_ld, const int dest_offset, __global real* dest, - const int upper, const int lower) { + const int upper, const int lower, + const int diagonal_imag_zero) { // Local memory to store a tile of the matrix (for coalescing) __local real tile[PADTRA_WPT*PADTRA_TILE][PADTRA_WPT*PADTRA_TILE + PADTRA_PAD]; @@ -147,6 +148,7 @@ __kernel void UnPadTransposeMatrix(const int src_one, const int src_two, // Stores the transposed value in the destination matrix if ((id_dest_one < dest_one) && (id_dest_two < dest_two)) { real value = tile[get_local_id(0)*PADTRA_WPT + w_two][get_local_id(1)*PADTRA_WPT + w_one]; + if (diagonal_imag_zero == 1 && id_dest_one == id_dest_two) { ImagToZero(value); } dest[id_dest_two*dest_ld + id_dest_one + dest_offset] = value; } } diff --git a/src/routine.cc b/src/routine.cc index 4b7ece41..d11edb0f 100644 --- a/src/routine.cc +++ b/src/routine.cc @@ -211,12 +211,13 @@ StatusCode Routine::PadCopyTransposeMatrix(const size_t src_one, const size_t sr const Buffer &dest, const bool do_transpose, const bool do_conjugate, const bool pad, const bool upper, const bool lower, + const bool diagonal_imag_zero, const Program &program) { // Determines whether or not the fast-version could potentially be used auto use_fast_kernel = (src_offset == 0) && (dest_offset == 0) && (do_conjugate == false) && (src_one == dest_one) && (src_two == dest_two) && (src_ld == dest_ld) && - (upper == false) && (lower == false); + (upper == false) && (lower == false) && (diagonal_imag_zero == false); // Determines the right kernel auto kernel_name = std::string{}; @@ -272,6 +273,7 @@ StatusCode Routine::PadCopyTransposeMatrix(const size_t src_one, const size_t sr else { kernel.SetArgument(10, static_cast<int>(upper)); kernel.SetArgument(11, static_cast<int>(lower)); + kernel.SetArgument(12, static_cast<int>(diagonal_imag_zero)); } } diff --git a/src/routines/xgemm.cc b/src/routines/xgemm.cc index 651ebb55..c8674282 100644 --- a/src/routines/xgemm.cc +++ b/src/routines/xgemm.cc @@ -108,18 +108,18 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout, // them up until they reach a certain multiple of size (kernel parameter dependent). status = PadCopyTransposeMatrix(a_one, a_two, a_ld, a_offset, a_buffer, m_ceiled, k_ceiled, m_ceiled, 0, temp_a, - a_do_transpose, a_conjugate, true, false, false, program); + a_do_transpose, a_conjugate, true, false, false, false, program); if (ErrorIn(status)) { return status; } status = PadCopyTransposeMatrix(b_one, b_two, b_ld, b_offset, b_buffer, n_ceiled, k_ceiled, n_ceiled, 0, temp_b, - b_do_transpose, b_conjugate, true, false, false, program); + b_do_transpose, b_conjugate, true, false, false, false, program); if (ErrorIn(status)) { return status; } // Only necessary for matrix C if it used both as input and output if (beta != static_cast<T>(0)) { status = PadCopyTransposeMatrix(c_one, c_two, c_ld, c_offset, c_buffer, m_ceiled, n_ceiled, m_ceiled, 0, temp_c, - c_do_transpose, false, true, false, false, program); + c_do_transpose, false, true, false, false, false, program); if (ErrorIn(status)) { return status; } } @@ -151,7 +151,7 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout, // Runs the post-processing kernel status = PadCopyTransposeMatrix(m_ceiled, n_ceiled, m_ceiled, 0, temp_c, c_one, c_two, c_ld, c_offset, c_buffer, - c_do_transpose, false, false, false, false, program); + c_do_transpose, false, false, false, false, false, program); if (ErrorIn(status)) { return status; } // Successfully finished the computation diff --git a/src/routines/xhemm.cc b/src/routines/xhemm.cc new file mode 100644 index 00000000..73f769ed --- /dev/null +++ b/src/routines/xhemm.cc @@ -0,0 +1,130 @@ + +// ================================================================================================= +// 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 Xhemm class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xhemm.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xhemm<T>::Xhemm(CommandQueue &queue, Event &event): + Xgemm<T>(queue, event) { +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xhemm<T>::DoHemm(const Layout layout, const Side side, const Triangle triangle, + const size_t m, const size_t n, + const T alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const Buffer &b_buffer, const size_t b_offset, const size_t b_ld, + const T beta, + const Buffer &c_buffer, const size_t c_offset, const size_t c_ld) { + + // Makes sure all dimensions are larger than zero + if ((m == 0) || (n == 0) ) { return StatusCode::kInvalidDimension; } + + // Computes the k dimension. This is based on whether or not the hermitian matrix is A (on the + // left) or B (on the right) in the Xgemm routine. + auto k = (side == Side::kLeft) ? m : n; + + // Checks for validity of the squared A matrix + auto status = TestMatrixA(k, k, a_buffer, a_offset, a_ld, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Determines which kernel to run based on the layout (the Xgemm kernel assumes column-major as + // default) and on whether we are dealing with an upper or lower triangle of the hermitian matrix + bool is_upper = ((triangle == Triangle::kUpper && layout != Layout::kRowMajor) || + (triangle == Triangle::kLower && layout == Layout::kRowMajor)); + auto kernel_name = (is_upper) ? "HermUpperToSquared" : "HermLowerToSquared"; + + // Temporary buffer for a copy of the hermitian matrix + try { + auto temp_herm = Buffer(context_, CL_MEM_READ_WRITE, k*k*sizeof(T)); + + // Creates a general matrix from the hermitian matrix to be able to run the regular Xgemm + // routine afterwards + try { + auto& program = GetProgramFromCache(); + auto kernel = Kernel(program, kernel_name); + + // Sets the arguments for the hermitian-to-squared kernel + kernel.SetArgument(0, static_cast<int>(k)); + kernel.SetArgument(1, static_cast<int>(a_ld)); + kernel.SetArgument(2, static_cast<int>(a_offset)); + kernel.SetArgument(3, a_buffer()); + kernel.SetArgument(4, static_cast<int>(k)); + kernel.SetArgument(5, static_cast<int>(k)); + kernel.SetArgument(6, static_cast<int>(0)); + kernel.SetArgument(7, temp_herm()); + + // Uses the common padding kernel's thread configuration. This is allowed, since the + // hermitian-to-squared kernel uses the same parameters. + auto global = std::vector<size_t>{Ceil(CeilDiv(k, db_["PAD_WPTX"]), db_["PAD_DIMX"]), + Ceil(CeilDiv(k, db_["PAD_WPTY"]), db_["PAD_DIMY"])}; + auto local = std::vector<size_t>{db_["PAD_DIMX"], db_["PAD_DIMY"]}; + status = RunKernel(kernel, global, local); + if (ErrorIn(status)) { return status; } + + // Runs the regular Xgemm code with either "C := AB+C" or ... + if (side == Side::kLeft) { + status = DoGemm(layout, Transpose::kNo, Transpose::kNo, + m, n, k, + alpha, + temp_herm, 0, k, + b_buffer, b_offset, b_ld, + beta, + c_buffer, c_offset, c_ld); + } + + // ... with "C := BA+C". Note that A and B are now reversed. + else { + status = DoGemm(layout, Transpose::kNo, Transpose::kNo, + m, n, k, + alpha, + b_buffer, b_offset, b_ld, + temp_herm, 0, k, + beta, + c_buffer, c_offset, c_ld); + + // A and B are now reversed, so also reverse the error codes returned from the Xgemm routine + switch(status) { + case StatusCode::kInvalidMatrixA: status = StatusCode::kInvalidMatrixB; break; + case StatusCode::kInvalidMatrixB: status = StatusCode::kInvalidMatrixA; break; + case StatusCode::kInvalidLeadDimA: status = StatusCode::kInvalidLeadDimB; break; + case StatusCode::kInvalidLeadDimB: status = StatusCode::kInvalidLeadDimA; break; + case StatusCode::kInsufficientMemoryA: status = StatusCode::kInsufficientMemoryB; break; + case StatusCode::kInsufficientMemoryB: status = StatusCode::kInsufficientMemoryA; break; + } + } + + // Return the status of the Xgemm routine + return status; + } catch (...) { return StatusCode::kInvalidKernel; } + } catch (...) { return StatusCode::kTempBufferAllocFailure; } +} + +// ================================================================================================= + +// Compiles the templated class +template class Xhemm<float2>; +template class Xhemm<double2>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/xher2k.cc b/src/routines/xher2k.cc new file mode 100644 index 00000000..b19b743b --- /dev/null +++ b/src/routines/xher2k.cc @@ -0,0 +1,178 @@ + +// ================================================================================================= +// 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 Xher2k class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xher2k.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xher2k<float2,float>::precision_ = Precision::kComplexSingle; +template <> const Precision Xher2k<double2,double>::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T, typename U> +Xher2k<T,U>::Xher2k(CommandQueue &queue, Event &event): + Routine(queue, event, {"Copy", "Pad", "Transpose", "PadTranspose", "Xgemm"}, precision_) { +} + +// ================================================================================================= + +// The main routine +template <typename T, typename U> +StatusCode Xher2k<T,U>::DoHer2k(const Layout layout, const Triangle triangle, const Transpose ab_transpose, + const size_t n, const size_t k, + const T alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const Buffer &b_buffer, const size_t b_offset, const size_t b_ld, + const U beta, + const Buffer &c_buffer, const size_t c_offset, const size_t c_ld) { + + // Makes sure all dimensions are larger than zero + if ((n == 0) || (k == 0) ) { return StatusCode::kInvalidDimension; } + + // Determines whether to apply the conjugate transpose to matrix B (argument: no transpose) or + // to matrix A (argument: conjugate transpose) + auto ab_conjugate = (ab_transpose != Transpose::kNo); + + // Computes whether or not the matrices are transposed in memory. This is based on their layout + // (row or column-major) and whether or not they are requested to be pre-transposed. + auto ab_rotated = (layout == Layout::kColMajor && ab_conjugate) || + (layout == Layout::kRowMajor && !ab_conjugate); + auto c_rotated = (layout == Layout::kRowMajor); + + // Computes the first and second dimensions of the A and B matrices taking the layout into account + auto ab_one = (ab_rotated) ? k : n; + auto ab_two = (ab_rotated) ? n : k; + + // Tests the matrices (A, B, C) for validity, first from a perspective of the OpenCL buffers and + // their sizes, and then from a perspective of parameter values (e.g. n, k). Tests whether the + // OpenCL buffers are valid and non-zero and whether the OpenCL buffers have sufficient storage + // space. Also tests that the leading dimensions of: + // matrix A cannot be less than N when rotated, or less than K when not-rotated + // matrix B cannot be less than N when rotated, or less than K when not-rotated + // matrix C cannot be less than N + auto status = TestMatrixA(ab_one, ab_two, a_buffer, a_offset, a_ld, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestMatrixB(ab_one, ab_two, b_buffer, b_offset, b_ld, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestMatrixC(n, n, c_buffer, c_offset, c_ld, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Calculates the ceiled versions of n and k + auto n_ceiled = Ceil(n, db_["NWG"]); + auto k_ceiled = Ceil(k, db_["KWG"]); + + // Decides which kernel to run: the upper-triangular or lower-triangular version + auto kernel_name = (triangle == Triangle::kUpper) ? "XgemmUpper" : "XgemmLower"; + + // Allocates space on the device for padded and/or transposed input and output matrices. + try { + auto temp_a1 = Buffer(context_, CL_MEM_READ_WRITE, k_ceiled*n_ceiled*sizeof(T)); + auto temp_b1 = Buffer(context_, CL_MEM_READ_WRITE, k_ceiled*n_ceiled*sizeof(T)); + auto temp_a2 = Buffer(context_, CL_MEM_READ_WRITE, k_ceiled*n_ceiled*sizeof(T)); + auto temp_b2 = Buffer(context_, CL_MEM_READ_WRITE, k_ceiled*n_ceiled*sizeof(T)); + auto temp_c = Buffer(context_, CL_MEM_READ_WRITE, n_ceiled*n_ceiled*sizeof(T)); + + // Loads the program from the database + auto& program = GetProgramFromCache(); + + // Runs the pre-processing kernels. This transposes the matrices A and B, but also pads zeros to + // fill them up until they reach a certain multiple of size (kernel parameter dependent). + status = PadCopyTransposeMatrix(ab_one, ab_two, a_ld, a_offset, a_buffer, + n_ceiled, k_ceiled, n_ceiled, 0, temp_a1, + ab_rotated, ab_conjugate, true, false, false, false, program); + if (ErrorIn(status)) { return status; } + status = PadCopyTransposeMatrix(ab_one, ab_two, a_ld, a_offset, a_buffer, + n_ceiled, k_ceiled, n_ceiled, 0, temp_a2, + ab_rotated, !ab_conjugate, true, false, false, false, program); + if (ErrorIn(status)) { return status; } + status = PadCopyTransposeMatrix(ab_one, ab_two, b_ld, b_offset, b_buffer, + n_ceiled, k_ceiled, n_ceiled, 0, temp_b1, + ab_rotated, ab_conjugate, true, false, false, false, program); + status = PadCopyTransposeMatrix(ab_one, ab_two, b_ld, b_offset, b_buffer, + n_ceiled, k_ceiled, n_ceiled, 0, temp_b2, + ab_rotated, !ab_conjugate, true, false, false, false, program); + if (ErrorIn(status)) { return status; } + + // Furthermore, also creates a (possibly padded) copy of matrix C, since it is not allowed to + // modify the other triangle. + status = PadCopyTransposeMatrix(n, n, c_ld, c_offset, c_buffer, + n_ceiled, n_ceiled, n_ceiled, 0, temp_c, + c_rotated, false, true, false, false, false, program); + if (ErrorIn(status)) { return status; } + + // Retrieves the XgemmUpper or XgemmLower kernel from the compiled binary + try { + auto kernel = Kernel(program, kernel_name); + + // Sets the kernel arguments + auto complex_beta = T{beta, static_cast<U>(0.0)}; + kernel.SetArgument(0, static_cast<int>(n_ceiled)); + kernel.SetArgument(1, static_cast<int>(k_ceiled)); + kernel.SetArgument(2, alpha); + kernel.SetArgument(3, complex_beta); + kernel.SetArgument(4, temp_a1()); + kernel.SetArgument(5, temp_b2()); + kernel.SetArgument(6, temp_c()); + + // Computes the global and local thread sizes + auto global = std::vector<size_t>{ + (n_ceiled * db_["MDIMC"]) / db_["MWG"], + (n_ceiled * db_["NDIMC"]) / db_["NWG"] + }; + auto local = std::vector<size_t>{db_["MDIMC"], db_["NDIMC"]}; + + // Launches the kernel + status = RunKernel(kernel, global, local); + if (ErrorIn(status)) { return status; } + + // Swaps the arguments for matrices A and B, sets 'beta' to 1, and conjugate alpha + auto conjugate_alpha = T{alpha.real(), -alpha.imag()}; + auto complex_one = T{static_cast<U>(1.0), static_cast<U>(0.0)}; + kernel.SetArgument(2, conjugate_alpha); + kernel.SetArgument(3, complex_one); + kernel.SetArgument(4, temp_b1()); + kernel.SetArgument(5, temp_a2()); + + // Runs the kernel again + status = RunKernel(kernel, global, local); + if (ErrorIn(status)) { return status; } + + // Runs the post-processing kernel + auto upper = (triangle == Triangle::kUpper); + auto lower = (triangle == Triangle::kLower); + status = PadCopyTransposeMatrix(n_ceiled, n_ceiled, n_ceiled, 0, temp_c, + n, n, c_ld, c_offset, c_buffer, + c_rotated, false, false, upper, lower, true, program); + if (ErrorIn(status)) { return status; } + + // Successfully finished the computation + return StatusCode::kSuccess; + } catch (...) { return StatusCode::kInvalidKernel; } + } catch (...) { return StatusCode::kTempBufferAllocFailure; } +} + +// ================================================================================================= + +// Compiles the templated class +template class Xher2k<float2,float>; +template class Xher2k<double2,double>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/xherk.cc b/src/routines/xherk.cc new file mode 100644 index 00000000..6bc9cd6c --- /dev/null +++ b/src/routines/xherk.cc @@ -0,0 +1,156 @@ + +// ================================================================================================= +// 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 Xherk class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xherk.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xherk<float2,float>::precision_ = Precision::kComplexSingle; +template <> const Precision Xherk<double2,double>::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T, typename U> +Xherk<T,U>::Xherk(CommandQueue &queue, Event &event): + Routine(queue, event, {"Copy", "Pad", "Transpose", "PadTranspose", "Xgemm"}, precision_) { +} + +// ================================================================================================= + +// The main routine +template <typename T, typename U> +StatusCode Xherk<T,U>::DoHerk(const Layout layout, const Triangle triangle, const Transpose a_transpose, + const size_t n, const size_t k, + const U alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const U beta, + const Buffer &c_buffer, const size_t c_offset, const size_t c_ld) { + + // Makes sure all dimensions are larger than zero + if ((n == 0) || (k == 0) ) { return StatusCode::kInvalidDimension; } + + // Determines whether to apply the conjugate transpose to matrix B (argument: no transpose) or + // to matrix A (argument: conjugate transpose) + auto a_conjugate = (a_transpose != Transpose::kNo); + auto b_conjugate = (a_transpose == Transpose::kNo); + + // Computes whether or not the matrices are transposed in memory. This is based on their layout + // (row or column-major) and whether or not they are requested to be pre-transposed. + auto a_rotated = (layout == Layout::kColMajor && a_conjugate) || + (layout == Layout::kRowMajor && !a_conjugate); + auto c_rotated = (layout == Layout::kRowMajor); + + // Computes the first and second dimensions of the A matrix taking the layout into account + auto a_one = (a_rotated) ? k : n; + auto a_two = (a_rotated) ? n : k; + + // Tests the two matrices (A, C) for validity, first from a perspective of the OpenCL buffers and + // their sizes, and then from a perspective of parameter values (e.g. n, k). Tests whether the + // OpenCL buffers are valid and non-zero and whether the OpenCL buffers have sufficient storage + // space. Also tests that the leading dimensions of: + // matrix A cannot be less than N when rotated, or less than K when not-rotated + // matrix C cannot be less than N + auto status = TestMatrixA(a_one, a_two, a_buffer, a_offset, a_ld, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestMatrixC(n, n, c_buffer, c_offset, c_ld, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Calculates the ceiled versions of n and k + auto n_ceiled = Ceil(n, db_["NWG"]); + auto k_ceiled = Ceil(k, db_["KWG"]); + + // Decides which kernel to run: the upper-triangular or lower-triangular version + auto kernel_name = (triangle == Triangle::kUpper) ? "XgemmUpper" : "XgemmLower"; + + // Allocates space on the device for padded and/or transposed input and output matrices. + try { + auto temp_a = Buffer(context_, CL_MEM_READ_WRITE, k_ceiled*n_ceiled*sizeof(T)); + auto temp_b = Buffer(context_, CL_MEM_READ_WRITE, k_ceiled*n_ceiled*sizeof(T)); + auto temp_c = Buffer(context_, CL_MEM_READ_WRITE, n_ceiled*n_ceiled*sizeof(T)); + + // Loads the program from the database + auto& program = GetProgramFromCache(); + + // Runs the pre-processing kernel. This transposes the matrix A, but also pads zeros to + // fill it up until it reaches a certain multiple of size (kernel parameter dependent). It + // creates two copies: + status = PadCopyTransposeMatrix(a_one, a_two, a_ld, a_offset, a_buffer, + n_ceiled, k_ceiled, n_ceiled, 0, temp_a, + a_rotated, a_conjugate, true, false, false, false, program); + if (ErrorIn(status)) { return status; } + status = PadCopyTransposeMatrix(a_one, a_two, a_ld, a_offset, a_buffer, + n_ceiled, k_ceiled, n_ceiled, 0, temp_b, + a_rotated, b_conjugate, true, false, false, false, program); + if (ErrorIn(status)) { return status; } + + // Furthermore, also creates a (possibly padded) copy of matrix C, since it is not allowed to + // modify the other triangle. + status = PadCopyTransposeMatrix(n, n, c_ld, c_offset, c_buffer, + n_ceiled, n_ceiled, n_ceiled, 0, temp_c, + c_rotated, false, true, false, false, false, program); + if (ErrorIn(status)) { return status; } + + // Retrieves the XgemmUpper or XgemmLower kernel from the compiled binary + try { + auto kernel = Kernel(program, kernel_name); + + // Sets the kernel arguments + auto complex_alpha = T{alpha, static_cast<U>(0.0)}; + auto complex_beta = T{beta, static_cast<U>(0.0)}; + kernel.SetArgument(0, static_cast<int>(n_ceiled)); + kernel.SetArgument(1, static_cast<int>(k_ceiled)); + kernel.SetArgument(2, complex_alpha); + kernel.SetArgument(3, complex_beta); + kernel.SetArgument(4, temp_a()); + kernel.SetArgument(5, temp_b()); + kernel.SetArgument(6, temp_c()); + + // Computes the global and local thread sizes + auto global = std::vector<size_t>{ + (n_ceiled * db_["MDIMC"]) / db_["MWG"], + (n_ceiled * db_["NDIMC"]) / db_["NWG"] + }; + auto local = std::vector<size_t>{db_["MDIMC"], db_["NDIMC"]}; + + // Launches the kernel + status = RunKernel(kernel, global, local); + if (ErrorIn(status)) { return status; } + + // Runs the post-processing kernel + auto upper = (triangle == Triangle::kUpper); + auto lower = (triangle == Triangle::kLower); + status = PadCopyTransposeMatrix(n_ceiled, n_ceiled, n_ceiled, 0, temp_c, + n, n, c_ld, c_offset, c_buffer, + c_rotated, false, false, upper, lower, true, program); + if (ErrorIn(status)) { return status; } + + // Successfully finished the computation + return StatusCode::kSuccess; + } catch (...) { return StatusCode::kInvalidKernel; } + } catch (...) { return StatusCode::kTempBufferAllocFailure; } +} + +// ================================================================================================= + +// Compiles the templated class +template class Xherk<float2,float>; +template class Xherk<double2,double>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/xsymm.cc b/src/routines/xsymm.cc index 97f35be8..b39eb24d 100644 --- a/src/routines/xsymm.cc +++ b/src/routines/xsymm.cc @@ -42,14 +42,14 @@ StatusCode Xsymm<T>::DoSymm(const Layout layout, const Side side, const Triangle // Computes the k dimension. This is based on whether or not the symmetric matrix is A (on the // left) or B (on the right) in the Xgemm routine. - size_t k = (side == Side::kLeft) ? m : n; + auto k = (side == Side::kLeft) ? m : n; // Checks for validity of the squared A matrix auto status = TestMatrixA(k, k, a_buffer, a_offset, a_ld, sizeof(T)); if (ErrorIn(status)) { return status; } // Determines which kernel to run based on the layout (the Xgemm kernel assumes column-major as - // default) and on whether we are dealing with an upper or lower triangle of the symmetrix matrix + // default) and on whether we are dealing with an upper or lower triangle of the symmetric matrix bool is_upper = ((triangle == Triangle::kUpper && layout != Layout::kRowMajor) || (triangle == Triangle::kLower && layout == Layout::kRowMajor)); auto kernel_name = (is_upper) ? "SymmUpperToSquared" : "SymmLowerToSquared"; @@ -75,7 +75,7 @@ StatusCode Xsymm<T>::DoSymm(const Layout layout, const Side side, const Triangle kernel.SetArgument(7, temp_symm()); // Uses the common padding kernel's thread configuration. This is allowed, since the - // symmetry-to-squared kernel uses the same parameters. + // symmetric-to-squared kernel uses the same parameters. auto global = std::vector<size_t>{Ceil(CeilDiv(k, db_["PAD_WPTX"]), db_["PAD_DIMX"]), Ceil(CeilDiv(k, db_["PAD_WPTY"]), db_["PAD_DIMY"])}; auto local = std::vector<size_t>{db_["PAD_DIMX"], db_["PAD_DIMY"]}; diff --git a/src/routines/xsyr2k.cc b/src/routines/xsyr2k.cc index a7aa6945..abb8b7eb 100644 --- a/src/routines/xsyr2k.cc +++ b/src/routines/xsyr2k.cc @@ -54,9 +54,6 @@ StatusCode Xsyr2k<T>::DoSyr2k(const Layout layout, const Triangle triangle, cons (layout == Layout::kRowMajor && ab_transpose == Transpose::kNo); auto c_rotated = (layout == Layout::kRowMajor); - // In case of complex data-types, the transpose can also become a conjugate transpose - auto ab_conjugate = (ab_transpose == Transpose::kConjugate); - // Computes the first and second dimensions of the A and B matrices taking the layout into account auto ab_one = (ab_rotated) ? k : n; auto ab_two = (ab_rotated) ? n : k; @@ -95,18 +92,18 @@ StatusCode Xsyr2k<T>::DoSyr2k(const Layout layout, const Triangle triangle, cons // fill them up until they reach a certain multiple of size (kernel parameter dependent). status = PadCopyTransposeMatrix(ab_one, ab_two, a_ld, a_offset, a_buffer, n_ceiled, k_ceiled, n_ceiled, 0, temp_a, - ab_rotated, ab_conjugate, true, false, false, program); + ab_rotated, false, true, false, false, false, program); if (ErrorIn(status)) { return status; } status = PadCopyTransposeMatrix(ab_one, ab_two, b_ld, b_offset, b_buffer, n_ceiled, k_ceiled, n_ceiled, 0, temp_b, - ab_rotated, ab_conjugate, true, false, false, program); + ab_rotated, false, true, false, false, false, program); if (ErrorIn(status)) { return status; } // Furthermore, also creates a (possibly padded) copy of matrix C, since it is not allowed to // modify the other triangle. status = PadCopyTransposeMatrix(n, n, c_ld, c_offset, c_buffer, n_ceiled, n_ceiled, n_ceiled, 0, temp_c, - c_rotated, false, true, false, false, program); + c_rotated, false, true, false, false, false, program); if (ErrorIn(status)) { return status; } // Retrieves the XgemmUpper or XgemmLower kernel from the compiled binary @@ -148,7 +145,7 @@ StatusCode Xsyr2k<T>::DoSyr2k(const Layout layout, const Triangle triangle, cons auto lower = (triangle == Triangle::kLower); status = PadCopyTransposeMatrix(n_ceiled, n_ceiled, n_ceiled, 0, temp_c, n, n, c_ld, c_offset, c_buffer, - c_rotated, false, false, upper, lower, program); + c_rotated, false, false, upper, lower, false, program); if (ErrorIn(status)) { return status; } // Successfully finished the computation diff --git a/src/routines/xsyrk.cc b/src/routines/xsyrk.cc index d8c150fd..3efa0598 100644 --- a/src/routines/xsyrk.cc +++ b/src/routines/xsyrk.cc @@ -53,9 +53,6 @@ StatusCode Xsyrk<T>::DoSyrk(const Layout layout, const Triangle triangle, const (layout == Layout::kRowMajor && a_transpose == Transpose::kNo); auto c_rotated = (layout == Layout::kRowMajor); - // In case of complex data-types, the transpose can also become a conjugate transpose - auto a_conjugate = (a_transpose == Transpose::kConjugate); - // Computes the first and second dimensions of the A matrix taking the layout into account auto a_one = (a_rotated) ? k : n; auto a_two = (a_rotated) ? n : k; @@ -87,17 +84,17 @@ StatusCode Xsyrk<T>::DoSyrk(const Layout layout, const Triangle triangle, const auto& program = GetProgramFromCache(); // Runs the pre-processing kernel. This transposes the matrix A, but also pads zeros to - // fill them up until they reach a certain multiple of size (kernel parameter dependent). + // fill it up until it reaches a certain multiple of size (kernel parameter dependent). status = PadCopyTransposeMatrix(a_one, a_two, a_ld, a_offset, a_buffer, n_ceiled, k_ceiled, n_ceiled, 0, temp_a, - a_rotated, a_conjugate, true, false, false, program); + a_rotated, false, true, false, false, false, program); if (ErrorIn(status)) { return status; } // Furthermore, also creates a (possibly padded) copy of matrix C, since it is not allowed to // modify the other triangle. status = PadCopyTransposeMatrix(n, n, c_ld, c_offset, c_buffer, n_ceiled, n_ceiled, n_ceiled, 0, temp_c, - c_rotated, false, true, false, false, program); + c_rotated, false, true, false, false, false, program); if (ErrorIn(status)) { return status; } // Retrieves the XgemmUpper or XgemmLower kernel from the compiled binary @@ -129,7 +126,7 @@ StatusCode Xsyrk<T>::DoSyrk(const Layout layout, const Triangle triangle, const auto lower = (triangle == Triangle::kLower); status = PadCopyTransposeMatrix(n_ceiled, n_ceiled, n_ceiled, 0, temp_c, n, n, c_ld, c_offset, c_buffer, - c_rotated, false, false, upper, lower, program); + c_rotated, false, false, upper, lower, false, program); if (ErrorIn(status)) { return status; } // Successfully finished the computation diff --git a/src/routines/xtrmm.cc b/src/routines/xtrmm.cc new file mode 100644 index 00000000..543df844 --- /dev/null +++ b/src/routines/xtrmm.cc @@ -0,0 +1,135 @@ + +// ================================================================================================= +// 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 Xtrmm class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xtrmm.h" + +#include <string> +#include <vector> + +namespace clblast { +// ================================================================================================= + +// Constructor: forwards to base class constructor +template <typename T> +Xtrmm<T>::Xtrmm(CommandQueue &queue, Event &event): + Xgemm<T>(queue, event) { +} + +// ================================================================================================= + +// The main routine +template <typename T> +StatusCode Xtrmm<T>::DoTrmm(const Layout layout, const Side side, const Triangle triangle, + const Transpose a_transpose, const Diagonal diagonal, + const size_t m, const size_t n, + const T alpha, + const Buffer &a_buffer, const size_t a_offset, const size_t a_ld, + const Buffer &b_buffer, const size_t b_offset, const size_t b_ld) { + + // Makes sure all dimensions are larger than zero + if ((m == 0) || (n == 0)) { return StatusCode::kInvalidDimension; } + + // Computes the k dimension. This is based on whether or not matrix is A (on the left) + // or B (on the right) in the Xgemm routine. + auto k = (side == Side::kLeft) ? m : n; + + // Checks for validity of the triangular A matrix + auto status = TestMatrixA(k, k, a_buffer, a_offset, a_ld, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Determines which kernel to run based on the layout (the Xgemm kernel assumes column-major as + // default) and on whether we are dealing with an upper or lower triangle of the triangular matrix + bool is_upper = ((triangle == Triangle::kUpper && layout != Layout::kRowMajor) || + (triangle == Triangle::kLower && layout == Layout::kRowMajor)); + auto kernel_name = (is_upper) ? "TrmmUpperToSquared" : "TrmmLowerToSquared"; + + // Determines whether or not the triangular matrix is unit-diagonal + auto unit_diagonal = (diagonal == Diagonal::kUnit) ? true : false; + + // Temporary buffer for a copy of the triangular matrix + try { + auto temp_triangular = Buffer(context_, CL_MEM_READ_WRITE, k*k*sizeof(T)); + + // Creates a general matrix from the triangular matrix to be able to run the regular Xgemm + // routine afterwards + try { + auto& program = GetProgramFromCache(); + auto kernel = Kernel(program, kernel_name); + + // Sets the arguments for the triangular-to-squared kernel + kernel.SetArgument(0, static_cast<int>(k)); + kernel.SetArgument(1, static_cast<int>(a_ld)); + kernel.SetArgument(2, static_cast<int>(a_offset)); + kernel.SetArgument(3, a_buffer()); + kernel.SetArgument(4, static_cast<int>(k)); + kernel.SetArgument(5, static_cast<int>(k)); + kernel.SetArgument(6, static_cast<int>(0)); + kernel.SetArgument(7, temp_triangular()); + kernel.SetArgument(8, static_cast<int>(unit_diagonal)); + + // Uses the common padding kernel's thread configuration. This is allowed, since the + // triangular-to-squared kernel uses the same parameters. + auto global = std::vector<size_t>{Ceil(CeilDiv(k, db_["PAD_WPTX"]), db_["PAD_DIMX"]), + Ceil(CeilDiv(k, db_["PAD_WPTY"]), db_["PAD_DIMY"])}; + auto local = std::vector<size_t>{db_["PAD_DIMX"], db_["PAD_DIMY"]}; + status = RunKernel(kernel, global, local); + if (ErrorIn(status)) { return status; } + + // Runs the regular Xgemm code with either "B := alpha*A*B" or ... + if (side == Side::kLeft) { + status = DoGemm(layout, a_transpose, Transpose::kNo, + m, n, k, + alpha, + temp_triangular, 0, k, + b_buffer, b_offset, b_ld, + static_cast<T>(0.0), + b_buffer, b_offset, b_ld); + } + + // ... with "B := alpha*B*A". Note that A and B are now reversed. + else { + status = DoGemm(layout, Transpose::kNo, a_transpose, + m, n, k, + alpha, + b_buffer, b_offset, b_ld, + temp_triangular, 0, k, + static_cast<T>(0.0), + b_buffer, b_offset, b_ld); + + // A and B are now reversed, so also reverse the error codes returned from the Xgemm routine + switch(status) { + case StatusCode::kInvalidMatrixA: status = StatusCode::kInvalidMatrixB; break; + case StatusCode::kInvalidMatrixB: status = StatusCode::kInvalidMatrixA; break; + case StatusCode::kInvalidLeadDimA: status = StatusCode::kInvalidLeadDimB; break; + case StatusCode::kInvalidLeadDimB: status = StatusCode::kInvalidLeadDimA; break; + case StatusCode::kInsufficientMemoryA: status = StatusCode::kInsufficientMemoryB; break; + case StatusCode::kInsufficientMemoryB: status = StatusCode::kInsufficientMemoryA; break; + } + } + + // Return the status of the Xgemm routine + return status; + } catch (...) { return StatusCode::kInvalidKernel; } + } catch (...) { return StatusCode::kTempBufferAllocFailure; } +} + +// ================================================================================================= + +// Compiles the templated class +template class Xtrmm<float>; +template class Xtrmm<double>; +template class Xtrmm<float2>; +template class Xtrmm<double2>; + +// ================================================================================================= +} // namespace clblast diff --git a/src/utilities.cc b/src/utilities.cc index 98570088..62abbb91 100644 --- a/src/utilities.cc +++ b/src/utilities.cc @@ -79,6 +79,13 @@ std::string ToString(Triangle value) { } } template <> +std::string ToString(Diagonal value) { + switch(value) { + case Diagonal::kUnit: return ToString(static_cast<int>(value))+" (unit)"; + case Diagonal::kNonUnit: return ToString(static_cast<int>(value))+" (non-unit)"; + } +} +template <> std::string ToString(Precision value) { switch(value) { case Precision::kHalf: return ToString(static_cast<int>(value))+" (half)"; @@ -143,6 +150,7 @@ template Layout GetArgument<Layout>(const int, char **, std::string&, const std: template Transpose GetArgument<Transpose>(const int, char **, std::string&, const std::string&, const Transpose); template Side GetArgument<Side>(const int, char **, std::string&, const std::string&, const Side); template Triangle GetArgument<Triangle>(const int, char **, std::string&, const std::string&, const Triangle); +template Diagonal GetArgument<Diagonal>(const int, char **, std::string&, const std::string&, const Diagonal); template Precision GetArgument<Precision>(const int, char **, std::string&, const std::string&, const Precision); // ================================================================================================= diff --git a/test/correctness/routines/xaxpy.cc b/test/correctness/routines/xaxpy.cc index 89315a0d..cf23ca9f 100644 --- a/test/correctness/routines/xaxpy.cc +++ b/test/correctness/routines/xaxpy.cc @@ -22,10 +22,10 @@ template <typename T> void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { // Creates a tester - TestBlas<T> tester{argc, argv, silent, name, TestXaxpy<T>::GetOptions(), - TestXaxpy<T>::RunRoutine, TestXaxpy<T>::RunReference, - TestXaxpy<T>::DownloadResult, TestXaxpy<T>::GetResultIndex, - TestXaxpy<T>::ResultID1, TestXaxpy<T>::ResultID2}; + TestBlas<T,T> tester{argc, argv, silent, name, TestXaxpy<T>::GetOptions(), + TestXaxpy<T>::RunRoutine, TestXaxpy<T>::RunReference, + TestXaxpy<T>::DownloadResult, TestXaxpy<T>::GetResultIndex, + TestXaxpy<T>::ResultID1, TestXaxpy<T>::ResultID2}; // This variable holds the arguments relevant for this routine auto args = Arguments<T>{}; diff --git a/test/correctness/routines/xgemm.cc b/test/correctness/routines/xgemm.cc index 72843d45..8a50e1ca 100644 --- a/test/correctness/routines/xgemm.cc +++ b/test/correctness/routines/xgemm.cc @@ -22,10 +22,10 @@ template <typename T> void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { // Creates a tester - TestBlas<T> tester{argc, argv, silent, name, TestXgemm<T>::GetOptions(), - TestXgemm<T>::RunRoutine, TestXgemm<T>::RunReference, - TestXgemm<T>::DownloadResult, TestXgemm<T>::GetResultIndex, - TestXgemm<T>::ResultID1, TestXgemm<T>::ResultID2}; + TestBlas<T,T> tester{argc, argv, silent, name, TestXgemm<T>::GetOptions(), + TestXgemm<T>::RunRoutine, TestXgemm<T>::RunReference, + TestXgemm<T>::DownloadResult, TestXgemm<T>::GetResultIndex, + TestXgemm<T>::ResultID1, TestXgemm<T>::ResultID2}; // This variable holds the arguments relevant for this routine auto args = Arguments<T>{}; diff --git a/test/correctness/routines/xgemv.cc b/test/correctness/routines/xgemv.cc index f1100810..50ce4699 100644 --- a/test/correctness/routines/xgemv.cc +++ b/test/correctness/routines/xgemv.cc @@ -22,10 +22,10 @@ template <typename T> void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { // Creates a tester - TestBlas<T> tester{argc, argv, silent, name, TestXgemv<T>::GetOptions(), - TestXgemv<T>::RunRoutine, TestXgemv<T>::RunReference, - TestXgemv<T>::DownloadResult, TestXgemv<T>::GetResultIndex, - TestXgemv<T>::ResultID1, TestXgemv<T>::ResultID2}; + TestBlas<T,T> tester{argc, argv, silent, name, TestXgemv<T>::GetOptions(), + TestXgemv<T>::RunRoutine, TestXgemv<T>::RunReference, + TestXgemv<T>::DownloadResult, TestXgemv<T>::GetResultIndex, + TestXgemv<T>::ResultID1, TestXgemv<T>::ResultID2}; // This variable holds the arguments relevant for this routine auto args = Arguments<T>{}; diff --git a/test/correctness/routines/xhemm.cc b/test/correctness/routines/xhemm.cc new file mode 100644 index 00000000..e8c82f65 --- /dev/null +++ b/test/correctness/routines/xhemm.cc @@ -0,0 +1,98 @@ + +// ================================================================================================= +// 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 Xhemm routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/xhemm.h" + +namespace clblast { +// ================================================================================================= + +// The correctness tester +template <typename T> +void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { + + // Creates a tester + TestBlas<T,T> tester{argc, argv, silent, name, TestXhemm<T>::GetOptions(), + TestXhemm<T>::RunRoutine, TestXhemm<T>::RunReference, + TestXhemm<T>::DownloadResult, TestXhemm<T>::GetResultIndex, + TestXhemm<T>::ResultID1, TestXhemm<T>::ResultID2}; + + // This variable holds the arguments relevant for this routine + auto args = Arguments<T>{}; + + // Loops over the test-cases from a data-layout point of view + for (auto &layout: tester.kLayouts) { args.layout = layout; + for (auto &side: tester.kSides) { args.side = side; + for (auto &triangle: tester.kTriangles) { args.triangle = triangle; + + // Creates the arguments vector for the regular tests + auto regular_test_vector = std::vector<Arguments<T>>{}; + for (auto &m: tester.kMatrixDims) { args.m = m; + for (auto &n: tester.kMatrixDims) { args.n = n; + for (auto &a_ld: tester.kMatrixDims) { args.a_ld = a_ld; + for (auto &a_offset: tester.kOffsets) { args.a_offset = a_offset; + for (auto &b_ld: tester.kMatrixDims) { args.b_ld = b_ld; + for (auto &b_offset: tester.kOffsets) { args.b_offset = b_offset; + for (auto &c_ld: tester.kMatrixDims) { args.c_ld = c_ld; + for (auto &c_offset: tester.kOffsets) { args.c_offset = c_offset; + for (auto &alpha: tester.kAlphaValues) { args.alpha = alpha; + for (auto &beta: tester.kBetaValues) { args.beta = beta; + args.a_size = TestXhemm<T>::GetSizeA(args); + args.b_size = TestXhemm<T>::GetSizeB(args); + args.c_size = TestXhemm<T>::GetSizeC(args); + if (args.a_size<1 || args.b_size<1 || args.c_size<1) { continue; } + regular_test_vector.push_back(args); + } + } + } + } + } + } + } + } + } + } + + // Creates the arguments vector for the invalid-buffer tests + auto invalid_test_vector = std::vector<Arguments<T>>{}; + args.m = args.n = tester.kBufferSize; + args.a_ld = args.b_ld = args.c_ld = tester.kBufferSize; + args.a_offset = args.b_offset = args.c_offset = 0; + for (auto &a_size: tester.kMatSizes) { args.a_size = a_size; + for (auto &b_size: tester.kMatSizes) { args.b_size = b_size; + for (auto &c_size: tester.kMatSizes) { args.c_size = c_size; + invalid_test_vector.push_back(args); + } + } + } + + // Runs the tests + const auto case_name = ToString(layout)+" "+ToString(side)+" "+ToString(triangle); + tester.TestRegular(regular_test_vector, case_name); + tester.TestInvalid(invalid_test_vector, case_name); + } + } + } +} + +// ================================================================================================= +} // namespace clblast + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::RunTest<clblast::float2>(argc, argv, true, "CHEMM"); + clblast::RunTest<clblast::double2>(argc, argv, true, "ZHEMM"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/routines/xher2k.cc b/test/correctness/routines/xher2k.cc new file mode 100644 index 00000000..7c0e5a92 --- /dev/null +++ b/test/correctness/routines/xher2k.cc @@ -0,0 +1,100 @@ + +// ================================================================================================= +// 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 Xher2k routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/xher2k.h" + +namespace clblast { +// ================================================================================================= + +// The correctness tester +template <typename T, typename U> +void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { + + // Creates a tester + TestBlas<T,U> tester{argc, argv, silent, name, TestXher2k<T,U>::GetOptions(), + TestXher2k<T,U>::RunRoutine, TestXher2k<T,U>::RunReference, + TestXher2k<T,U>::DownloadResult, TestXher2k<T,U>::GetResultIndex, + TestXher2k<T,U>::ResultID1, TestXher2k<T,U>::ResultID2}; + + // This variable holds the arguments relevant for this routine + auto args = Arguments<U>{}; + + // Loops over the test-cases from a data-layout point of view + for (auto &layout: tester.kLayouts) { args.layout = layout; + for (auto &triangle: tester.kTriangles) { args.triangle = triangle; + for (auto &ab_transpose: {Transpose::kNo, Transpose::kConjugate}) { // Regular transpose not a + args.a_transpose = ab_transpose; // valid BLAS option + args.b_transpose = ab_transpose; + + // Creates the arguments vector for the regular tests + auto regular_test_vector = std::vector<Arguments<U>>{}; + for (auto &n: tester.kMatrixDims) { args.n = n; + for (auto &k: tester.kMatrixDims) { args.k = k; + for (auto &a_ld: tester.kMatrixDims) { args.a_ld = a_ld; + for (auto &a_offset: tester.kOffsets) { args.a_offset = a_offset; + for (auto &b_ld: tester.kMatrixDims) { args.b_ld = b_ld; + for (auto &b_offset: tester.kOffsets) { args.b_offset = b_offset; + for (auto &c_ld: tester.kMatrixDims) { args.c_ld = c_ld; + for (auto &c_offset: tester.kOffsets) { args.c_offset = c_offset; + for (auto &alpha: tester.kAlphaValues) { args.alpha = alpha; + for (auto &beta: tester.kBetaValues) { args.beta = beta; + args.a_size = TestXher2k<T,U>::GetSizeA(args); + args.b_size = TestXher2k<T,U>::GetSizeB(args); + args.c_size = TestXher2k<T,U>::GetSizeC(args); + if (args.a_size<1 || args.b_size<1 || args.c_size<1) { continue; } + regular_test_vector.push_back(args); + } + } + } + } + } + } + } + } + } + } + + // Creates the arguments vector for the invalid-buffer tests + auto invalid_test_vector = std::vector<Arguments<U>>{}; + args.n = args.k = tester.kBufferSize; + args.a_ld = args.b_ld = args.c_ld = tester.kBufferSize; + args.a_offset = args.b_offset = args.c_offset = 0; + for (auto &a_size: tester.kMatSizes) { args.a_size = a_size; + for (auto &b_size: tester.kMatSizes) { args.b_size = b_size; + for (auto &c_size: tester.kMatSizes) { args.c_size = c_size; + invalid_test_vector.push_back(args); + } + } + } + + // Runs the tests + const auto case_name = ToString(layout)+" "+ToString(triangle)+" "+ToString(ab_transpose); + tester.TestRegular(regular_test_vector, case_name); + tester.TestInvalid(invalid_test_vector, case_name); + } + } + } +} + +// ================================================================================================= +} // namespace clblast + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::RunTest<clblast::float2,float>(argc, argv, false, "CHER2K"); + clblast::RunTest<clblast::double2,double>(argc, argv, true, "ZHER2K"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/routines/xherk.cc b/test/correctness/routines/xherk.cc new file mode 100644 index 00000000..dc5c6caf --- /dev/null +++ b/test/correctness/routines/xherk.cc @@ -0,0 +1,92 @@ + +// ================================================================================================= +// 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 Xherk routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/xherk.h" + +namespace clblast { +// ================================================================================================= + +// The correctness tester +template <typename T, typename U> +void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { + + // Creates a tester + TestBlas<T,U> tester{argc, argv, silent, name, TestXherk<T,U>::GetOptions(), + TestXherk<T,U>::RunRoutine, TestXherk<T,U>::RunReference, + TestXherk<T,U>::DownloadResult, TestXherk<T,U>::GetResultIndex, + TestXherk<T,U>::ResultID1, TestXherk<T,U>::ResultID2}; + + // This variable holds the arguments relevant for this routine + auto args = Arguments<U>{}; + + // Loops over the test-cases from a data-layout point of view + for (auto &layout: tester.kLayouts) { args.layout = layout; + for (auto &triangle: tester.kTriangles) { args.triangle = triangle; + for (auto &a_transpose: {Transpose::kNo, Transpose::kConjugate}) { // Regular transpose not a + args.a_transpose = a_transpose; // valid BLAS option + + // Creates the arguments vector for the regular tests + auto regular_test_vector = std::vector<Arguments<U>>{}; + for (auto &n: tester.kMatrixDims) { args.n = n; + for (auto &k: tester.kMatrixDims) { args.k = k; + for (auto &a_ld: tester.kMatrixDims) { args.a_ld = a_ld; + for (auto &a_offset: tester.kOffsets) { args.a_offset = a_offset; + for (auto &c_ld: tester.kMatrixDims) { args.c_ld = c_ld; + for (auto &c_offset: tester.kOffsets) { args.c_offset = c_offset; + for (auto &alpha: tester.kAlphaValues) { args.alpha = alpha; + for (auto &beta: tester.kBetaValues) { args.beta = beta; + args.a_size = TestXherk<T,U>::GetSizeA(args); + args.c_size = TestXherk<T,U>::GetSizeC(args); + if (args.a_size<1 || args.c_size<1) { continue; } + regular_test_vector.push_back(args); + } + } + } + } + } + } + } + } + + // Creates the arguments vector for the invalid-buffer tests + auto invalid_test_vector = std::vector<Arguments<U>>{}; + args.n = args.k = tester.kBufferSize; + args.a_ld = args.c_ld = tester.kBufferSize; + args.a_offset = args.c_offset = 0; + for (auto &a_size: tester.kMatSizes) { args.a_size = a_size; + for (auto &c_size: tester.kMatSizes) { args.c_size = c_size; + invalid_test_vector.push_back(args); + } + } + + // Runs the tests + const auto case_name = ToString(layout)+" "+ToString(triangle)+" "+ToString(a_transpose); + tester.TestRegular(regular_test_vector, case_name); + tester.TestInvalid(invalid_test_vector, case_name); + } + } + } +} + +// ================================================================================================= +} // namespace clblast + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::RunTest<clblast::float2,float>(argc, argv, false, "CHERK"); + clblast::RunTest<clblast::double2,double>(argc, argv, true, "ZHERK"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/routines/xsymm.cc b/test/correctness/routines/xsymm.cc index 3da654c3..a919a056 100644 --- a/test/correctness/routines/xsymm.cc +++ b/test/correctness/routines/xsymm.cc @@ -22,10 +22,10 @@ template <typename T> void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { // Creates a tester - TestBlas<T> tester{argc, argv, silent, name, TestXsymm<T>::GetOptions(), - TestXsymm<T>::RunRoutine, TestXsymm<T>::RunReference, - TestXsymm<T>::DownloadResult, TestXsymm<T>::GetResultIndex, - TestXsymm<T>::ResultID1, TestXsymm<T>::ResultID2}; + TestBlas<T,T> tester{argc, argv, silent, name, TestXsymm<T>::GetOptions(), + TestXsymm<T>::RunRoutine, TestXsymm<T>::RunReference, + TestXsymm<T>::DownloadResult, TestXsymm<T>::GetResultIndex, + TestXsymm<T>::ResultID1, TestXsymm<T>::ResultID2}; // This variable holds the arguments relevant for this routine auto args = Arguments<T>{}; diff --git a/test/correctness/routines/xsyr2k.cc b/test/correctness/routines/xsyr2k.cc index 8b03087c..736aa4e5 100644 --- a/test/correctness/routines/xsyr2k.cc +++ b/test/correctness/routines/xsyr2k.cc @@ -22,10 +22,10 @@ template <typename T> void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { // Creates a tester - TestBlas<T> tester{argc, argv, silent, name, TestXsyr2k<T>::GetOptions(), - TestXsyr2k<T>::RunRoutine, TestXsyr2k<T>::RunReference, - TestXsyr2k<T>::DownloadResult, TestXsyr2k<T>::GetResultIndex, - TestXsyr2k<T>::ResultID1, TestXsyr2k<T>::ResultID2}; + TestBlas<T,T> tester{argc, argv, silent, name, TestXsyr2k<T>::GetOptions(), + TestXsyr2k<T>::RunRoutine, TestXsyr2k<T>::RunReference, + TestXsyr2k<T>::DownloadResult, TestXsyr2k<T>::GetResultIndex, + TestXsyr2k<T>::ResultID1, TestXsyr2k<T>::ResultID2}; // This variable holds the arguments relevant for this routine auto args = Arguments<T>{}; diff --git a/test/correctness/routines/xsyrk.cc b/test/correctness/routines/xsyrk.cc index d4552a78..a62a0ebf 100644 --- a/test/correctness/routines/xsyrk.cc +++ b/test/correctness/routines/xsyrk.cc @@ -22,10 +22,10 @@ template <typename T> void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { // Creates a tester - TestBlas<T> tester{argc, argv, silent, name, TestXsyrk<T>::GetOptions(), - TestXsyrk<T>::RunRoutine, TestXsyrk<T>::RunReference, - TestXsyrk<T>::DownloadResult, TestXsyrk<T>::GetResultIndex, - TestXsyrk<T>::ResultID1, TestXsyrk<T>::ResultID2}; + TestBlas<T,T> tester{argc, argv, silent, name, TestXsyrk<T>::GetOptions(), + TestXsyrk<T>::RunRoutine, TestXsyrk<T>::RunReference, + TestXsyrk<T>::DownloadResult, TestXsyrk<T>::GetResultIndex, + TestXsyrk<T>::ResultID1, TestXsyrk<T>::ResultID2}; // This variable holds the arguments relevant for this routine auto args = Arguments<T>{}; diff --git a/test/correctness/routines/xtrmm.cc b/test/correctness/routines/xtrmm.cc new file mode 100644 index 00000000..0bb6294c --- /dev/null +++ b/test/correctness/routines/xtrmm.cc @@ -0,0 +1,96 @@ + +// ================================================================================================= +// 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 Xtrmm routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/xtrmm.h" + +namespace clblast { +// ================================================================================================= + +// The correctness tester +template <typename T> +void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { + + // Creates a tester + TestBlas<T,T> tester{argc, argv, silent, name, TestXtrmm<T>::GetOptions(), + TestXtrmm<T>::RunRoutine, TestXtrmm<T>::RunReference, + TestXtrmm<T>::DownloadResult, TestXtrmm<T>::GetResultIndex, + TestXtrmm<T>::ResultID1, TestXtrmm<T>::ResultID2}; + + // This variable holds the arguments relevant for this routine + auto args = Arguments<T>{}; + + // Loops over the test-cases from a data-layout point of view + for (auto &layout: tester.kLayouts) { args.layout = layout; + for (auto &side: tester.kSides) { args.side = side; + for (auto &triangle: tester.kTriangles) { args.triangle = triangle; + for (auto &a_transpose: tester.kTransposes) { args.a_transpose = a_transpose; + for (auto &diagonal: tester.kDiagonals) { args.diagonal = diagonal; + + // Creates the arguments vector for the regular tests + auto regular_test_vector = std::vector<Arguments<T>>{}; + for (auto &m: tester.kMatrixDims) { args.m = m; + for (auto &n: tester.kMatrixDims) { args.n = n; + for (auto &a_ld: tester.kMatrixDims) { args.a_ld = a_ld; + for (auto &a_offset: tester.kOffsets) { args.a_offset = a_offset; + for (auto &b_ld: tester.kMatrixDims) { args.b_ld = b_ld; + for (auto &b_offset: tester.kOffsets) { args.b_offset = b_offset; + for (auto &alpha: tester.kAlphaValues) { args.alpha = alpha; + args.a_size = TestXtrmm<T>::GetSizeA(args); + args.b_size = TestXtrmm<T>::GetSizeB(args); + if (args.a_size<1 || args.b_size<1) { continue; } + regular_test_vector.push_back(args); + } + } + } + } + } + } + } + + // Creates the arguments vector for the invalid-buffer tests + auto invalid_test_vector = std::vector<Arguments<T>>{}; + args.m = args.n = tester.kBufferSize; + args.a_ld = args.b_ld = tester.kBufferSize; + args.a_offset = args.b_offset = 0; + for (auto &a_size: tester.kMatSizes) { args.a_size = a_size; + for (auto &b_size: tester.kMatSizes) { args.b_size = b_size; + invalid_test_vector.push_back(args); + } + } + + // Runs the tests + const auto case_name = ToString(layout)+" "+ToString(side)+" "+ToString(triangle)+" "+ + ToString(a_transpose)+" "+ToString(diagonal); + tester.TestRegular(regular_test_vector, case_name); + tester.TestInvalid(invalid_test_vector, case_name); + } + } + } + } + } +} + +// ================================================================================================= +} // namespace clblast + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + clblast::RunTest<float>(argc, argv, false, "STRMM"); + clblast::RunTest<double>(argc, argv, true, "DTRMM"); + clblast::RunTest<clblast::float2>(argc, argv, true, "CTRMM"); + clblast::RunTest<clblast::double2>(argc, argv, true, "ZTRMM"); + return 0; +} + +// ================================================================================================= diff --git a/test/correctness/testblas.cc b/test/correctness/testblas.cc index 0e72e429..5951b177 100644 --- a/test/correctness/testblas.cc +++ b/test/correctness/testblas.cc @@ -19,21 +19,23 @@ namespace clblast { // ================================================================================================= // The transpose-options to test with (data-type dependent) -template <> const std::vector<Transpose> TestBlas<float>::kTransposes = {Transpose::kNo, Transpose::kYes}; -template <> const std::vector<Transpose> TestBlas<double>::kTransposes = {Transpose::kNo, Transpose::kYes}; -template <> const std::vector<Transpose> TestBlas<float2>::kTransposes = {Transpose::kNo, Transpose::kYes, Transpose::kConjugate}; -template <> const std::vector<Transpose> TestBlas<double2>::kTransposes = {Transpose::kNo, Transpose::kYes, Transpose::kConjugate}; +template <> const std::vector<Transpose> TestBlas<float,float>::kTransposes = {Transpose::kNo, Transpose::kYes}; +template <> const std::vector<Transpose> TestBlas<double,double>::kTransposes = {Transpose::kNo, Transpose::kYes}; +template <> const std::vector<Transpose> TestBlas<float2,float2>::kTransposes = {Transpose::kNo, Transpose::kYes, Transpose::kConjugate}; +template <> const std::vector<Transpose> TestBlas<double2,double2>::kTransposes = {Transpose::kNo, Transpose::kYes, Transpose::kConjugate}; +template <> const std::vector<Transpose> TestBlas<float2,float>::kTransposes = {Transpose::kNo, Transpose::kConjugate}; +template <> const std::vector<Transpose> TestBlas<double2,double>::kTransposes = {Transpose::kNo, Transpose::kConjugate}; // ================================================================================================= // Constructor, initializes the base class tester and input data -template <typename T> -TestBlas<T>::TestBlas(int argc, char *argv[], const bool silent, - const std::string &name, const std::vector<std::string> &options, - const Routine run_routine, const Routine run_reference, - const ResultGet get_result, const ResultIndex get_index, - const ResultIterator get_id1, const ResultIterator get_id2): - Tester<T>{argc, argv, silent, name, options}, +template <typename T, typename U> +TestBlas<T,U>::TestBlas(int argc, char *argv[], const bool silent, + const std::string &name, const std::vector<std::string> &options, + const Routine run_routine, const Routine run_reference, + const ResultGet get_result, const ResultIndex get_index, + const ResultIterator get_id1, const ResultIterator get_id2): + Tester<T,U>{argc, argv, silent, name, options}, run_routine_(run_routine), run_reference_(run_reference), get_result_(get_result), @@ -65,9 +67,9 @@ TestBlas<T>::TestBlas(int argc, char *argv[], const bool silent, // =============================================================================================== // Tests the routine for a wide variety of parameters -template <typename T> -void TestBlas<T>::TestRegular(std::vector<Arguments<T>> &test_vector, const std::string &name) { - if (!PrecisionSupported()) { return; } +template <typename T, typename U> +void TestBlas<T,U>::TestRegular(std::vector<Arguments<U>> &test_vector, const std::string &name) { + if (!PrecisionSupported<T>(device_)) { return; } TestStart("regular behaviour", name); // Iterates over all the to-be-tested combinations of arguments @@ -132,9 +134,9 @@ void TestBlas<T>::TestRegular(std::vector<Arguments<T>> &test_vector, const std: // Tests the routine for cases with invalid OpenCL memory buffer sizes. Tests only on return-types, // does not test for results (if any). -template <typename T> -void TestBlas<T>::TestInvalid(std::vector<Arguments<T>> &test_vector, const std::string &name) { - if (!PrecisionSupported()) { return; } +template <typename T, typename U> +void TestBlas<T,U>::TestInvalid(std::vector<Arguments<U>> &test_vector, const std::string &name) { + if (!PrecisionSupported<T>(device_)) { return; } TestStart("invalid buffer sizes", name); // Iterates over all the to-be-tested combinations of arguments @@ -176,10 +178,12 @@ void TestBlas<T>::TestInvalid(std::vector<Arguments<T>> &test_vector, const std: // ================================================================================================= // Compiles the templated class -template class TestBlas<float>; -template class TestBlas<double>; -template class TestBlas<float2>; -template class TestBlas<double2>; +template class TestBlas<float, float>; +template class TestBlas<double, double>; +template class TestBlas<float2, float2>; +template class TestBlas<double2, double2>; +template class TestBlas<float2, float>; +template class TestBlas<double2, double>; // ================================================================================================= } // namespace clblast diff --git a/test/correctness/testblas.h b/test/correctness/testblas.h index 1f92cb30..96c140c1 100644 --- a/test/correctness/testblas.h +++ b/test/correctness/testblas.h @@ -9,6 +9,8 @@ // // This file tests any CLBlast routine. It contains two types of tests: one testing all sorts of // input combinations, and one deliberatly testing with invalid values. +// Typename T: the data-type of the routine's memory buffers (==precision) +// Typename U: the data-type of the alpha and beta arguments // // ================================================================================================= @@ -24,23 +26,22 @@ namespace clblast { // ================================================================================================= // See comment at top of file for a description of the class -template <typename T> -class TestBlas: public Tester<T> { +template <typename T, typename U> +class TestBlas: public Tester<T,U> { public: // Uses several variables from the Tester class - using Tester<T>::context_; - using Tester<T>::queue_; + using Tester<T,U>::context_; + using Tester<T,U>::queue_; + using Tester<T,U>::full_test_; + using Tester<T,U>::device_; // Uses several helper functions from the Tester class - using Tester<T>::TestStart; - using Tester<T>::TestEnd; - using Tester<T>::TestSimilarity; - using Tester<T>::TestErrorCount; - using Tester<T>::TestErrorCodes; - using Tester<T>::GetExampleScalars; - using Tester<T>::GetOffsets; - using Tester<T>::PrecisionSupported; + using Tester<T,U>::TestStart; + using Tester<T,U>::TestEnd; + using Tester<T,U>::TestErrorCount; + using Tester<T,U>::TestErrorCodes; + using Tester<T,U>::GetOffsets; // Test settings for the regular test. Append to these lists in case more tests are required. const std::vector<size_t> kVectorDims = { 7, 93, 4096 }; @@ -48,8 +49,8 @@ class TestBlas: public Tester<T> { const std::vector<size_t> kMatrixDims = { 7, 64 }; const std::vector<size_t> kMatrixVectorDims = { 61, 512 }; const std::vector<size_t> kOffsets = GetOffsets(); - const std::vector<T> kAlphaValues = GetExampleScalars(); - const std::vector<T> kBetaValues = GetExampleScalars(); + const std::vector<U> kAlphaValues = GetExampleScalars<U>(full_test_); + const std::vector<U> kBetaValues = GetExampleScalars<U>(full_test_); // Test settings for the invalid tests const std::vector<size_t> kInvalidIncrements = { 0, 1 }; @@ -61,13 +62,14 @@ class TestBlas: public Tester<T> { const std::vector<Layout> kLayouts = {Layout::kRowMajor, Layout::kColMajor}; const std::vector<Triangle> kTriangles = {Triangle::kUpper, Triangle::kLower}; const std::vector<Side> kSides = {Side::kLeft, Side::kRight}; + const std::vector<Diagonal> kDiagonals = {Diagonal::kUnit, Diagonal::kNonUnit}; static const std::vector<Transpose> kTransposes; // Data-type dependent, see .cc-file // Shorthand for the routine-specific functions passed to the tester - using Routine = std::function<StatusCode(const Arguments<T>&, const Buffers&, CommandQueue&)>; - using ResultGet = std::function<std::vector<T>(const Arguments<T>&, Buffers&, CommandQueue&)>; - using ResultIndex = std::function<size_t(const Arguments<T>&, const size_t, const size_t)>; - using ResultIterator = std::function<size_t(const Arguments<T>&)>; + using Routine = std::function<StatusCode(const Arguments<U>&, const Buffers&, CommandQueue&)>; + using ResultGet = std::function<std::vector<T>(const Arguments<U>&, Buffers&, CommandQueue&)>; + using ResultIndex = std::function<size_t(const Arguments<U>&, const size_t, const size_t)>; + using ResultIterator = std::function<size_t(const Arguments<U>&)>; // Constructor, initializes the base class tester and input data TestBlas(int argc, char *argv[], const bool silent, @@ -76,8 +78,8 @@ class TestBlas: public Tester<T> { const ResultIndex get_index, const ResultIterator get_id1, const ResultIterator get_id2); // The test functions, taking no inputs - void TestRegular(std::vector<Arguments<T>> &test_vector, const std::string &name); - void TestInvalid(std::vector<Arguments<T>> &test_vector, const std::string &name); + void TestRegular(std::vector<Arguments<U>> &test_vector, const std::string &name); + void TestInvalid(std::vector<Arguments<U>> &test_vector, const std::string &name); private: diff --git a/test/correctness/tester.cc b/test/correctness/tester.cc index 4a179718..378968ed 100644 --- a/test/correctness/tester.cc +++ b/test/correctness/tester.cc @@ -23,9 +23,9 @@ namespace clblast { // General constructor for all CLBlast testers. It prints out the test header to stdout and sets-up // the clBLAS library for reference. -template <typename T> -Tester<T>::Tester(int argc, char *argv[], const bool silent, - const std::string &name, const std::vector<std::string> &options): +template <typename T, typename U> +Tester<T,U>::Tester(int argc, char *argv[], const bool silent, + const std::string &name, const std::vector<std::string> &options): help_("Options given/available:\n"), platform_(Platform(GetArgument(argc, argv, help_, kArgPlatform, size_t{0}))), device_(Device(platform_, kDeviceType, GetArgument(argc, argv, help_, kArgDevice, size_t{0}))), @@ -51,7 +51,7 @@ Tester<T>::Tester(int argc, char *argv[], const bool silent, kPrintMessage.c_str(), name.c_str(), kPrintEnd.c_str()); // Checks whether the precision is supported - if (!PrecisionSupported()) { + if (!PrecisionSupported<T>(device_)) { fprintf(stdout, "\n* All tests skipped: %sUnsupported precision%s\n", kPrintWarning.c_str(), kPrintEnd.c_str()); return; @@ -76,9 +76,9 @@ Tester<T>::Tester(int argc, char *argv[], const bool silent, } // Destructor prints the summary of the test cases and cleans-up the clBLAS library -template <typename T> -Tester<T>::~Tester() { - if (PrecisionSupported()) { +template <typename T, typename U> +Tester<T,U>::~Tester() { + if (PrecisionSupported<T>(device_)) { fprintf(stdout, "* Completed all test-cases for this routine. Results:\n"); fprintf(stdout, " %lu test(s) passed\n", tests_passed_); if (tests_skipped_ > 0) { fprintf(stdout, "%s", kPrintWarning.c_str()); } @@ -94,8 +94,8 @@ Tester<T>::~Tester() { // Function called at the start of each test. This prints a header with information about the // test and re-initializes all test data-structures. -template <typename T> -void Tester<T>::TestStart(const std::string &test_name, const std::string &test_configuration) { +template <typename T, typename U> +void Tester<T,U>::TestStart(const std::string &test_name, const std::string &test_configuration) { // Prints the header fprintf(stdout, "* Testing %s'%s'%s for %s'%s'%s:\n", @@ -113,8 +113,8 @@ void Tester<T>::TestStart(const std::string &test_name, const std::string &test_ // Function called at the end of each test. This prints errors if any occured. It also prints a // summary of the number of sub-tests passed/failed. -template <typename T> -void Tester<T>::TestEnd() { +template <typename T, typename U> +void Tester<T,U>::TestEnd() { fprintf(stdout, "\n"); tests_passed_ += num_passed_; tests_failed_ += num_skipped_; @@ -137,6 +137,7 @@ void Tester<T>::TestEnd() { if (o == kArgBTransp) { fprintf(stdout, "%s=%d ", kArgBTransp, entry.args.b_transpose);} if (o == kArgSide) { fprintf(stdout, "%s=%d ", kArgSide, entry.args.side);} if (o == kArgTriangle) { fprintf(stdout, "%s=%d ", kArgTriangle, entry.args.triangle);} + if (o == kArgDiagonal) { fprintf(stdout, "%s=%d ", kArgDiagonal, entry.args.diagonal);} if (o == kArgXInc) { fprintf(stdout, "%s=%lu ", kArgXInc, entry.args.x_inc);} if (o == kArgYInc) { fprintf(stdout, "%s=%lu ", kArgYInc, entry.args.y_inc);} if (o == kArgXOffset) { fprintf(stdout, "%s=%lu ", kArgXOffset, entry.args.x_offset);} @@ -171,45 +172,9 @@ void Tester<T>::TestEnd() { // ================================================================================================= -// Compares two floating point values and returns whether they are within an acceptable error -// margin. This replaces GTest's EXPECT_NEAR(). -template <typename T> -bool Tester<T>::TestSimilarity(const T val1, const T val2) { - const auto difference = std::fabs(val1 - val2); - - // Shortcut, handles infinities - if (val1 == val2) { - return true; - } - // The values are zero or very small: the relative error is less meaningful - else if (val1 == 0 || val2 == 0 || difference < static_cast<T>(kErrorMarginAbsolute)) { - return (difference < static_cast<T>(kErrorMarginAbsolute)); - } - // Use relative error - else { - return (difference / (std::fabs(val1)+std::fabs(val2))) < static_cast<T>(kErrorMarginRelative); - } -} - -// Specialisations for complex data-types -template <> -bool Tester<float2>::TestSimilarity(const float2 val1, const float2 val2) { - auto real = Tester<float>::TestSimilarity(val1.real(), val2.real()); - auto imag = Tester<float>::TestSimilarity(val1.imag(), val2.imag()); - return (real && imag); -} -template <> -bool Tester<double2>::TestSimilarity(const double2 val1, const double2 val2) { - auto real = Tester<double>::TestSimilarity(val1.real(), val2.real()); - auto imag = Tester<double>::TestSimilarity(val1.imag(), val2.imag()); - return (real && imag); -} - -// ================================================================================================= - // Handles a 'pass' or 'error' depending on whether there are any errors -template <typename T> -void Tester<T>::TestErrorCount(const size_t errors, const size_t size, const Arguments<T> &args) { +template <typename T, typename U> +void Tester<T,U>::TestErrorCount(const size_t errors, const size_t size, const Arguments<U> &args) { // Finished successfully if (errors == 0) { @@ -227,9 +192,9 @@ void Tester<T>::TestErrorCount(const size_t errors, const size_t size, const Arg // Compares two status codes for equality. The outcome can be a pass (they are the same), a warning // (CLBlast reported a compilation error), or an error (they are different). -template <typename T> -void Tester<T>::TestErrorCodes(const StatusCode clblas_status, const StatusCode clblast_status, - const Arguments<T> &args) { +template <typename T, typename U> +void Tester<T,U>::TestErrorCodes(const StatusCode clblas_status, const StatusCode clblast_status, + const Arguments<U> &args) { // Finished successfully if (clblas_status == clblast_status) { @@ -260,62 +225,26 @@ void Tester<T>::TestErrorCodes(const StatusCode clblas_status, const StatusCode // ================================================================================================= -// Retrieves a list of example scalar values, used for the alpha and beta arguments for the various -// routines. This function is specialised for the different data-types. -template <> -const std::vector<float> Tester<float>::GetExampleScalars() { - if (full_test_) { return {0.0f, 1.0f, 3.14f}; } - else { return {3.14f}; } -} -template <> -const std::vector<double> Tester<double>::GetExampleScalars() { - if (full_test_) { return {0.0, 1.0, 3.14}; } - else { return {3.14}; } -} -template <> -const std::vector<float2> Tester<float2>::GetExampleScalars() { - if (full_test_) { return {{0.0f, 0.0f}, {1.0f, 1.3f}, {2.42f, 3.14f}}; } - else { return {{2.42f, 3.14f}}; } -} -template <> -const std::vector<double2> Tester<double2>::GetExampleScalars() { - if (full_test_) { return {{0.0, 0.0}, {1.0, 1.3}, {2.42, 3.14}}; } - else { return {{2.42, 3.14}}; } -} - // Retrieves the offset values to test with -template <typename T> -const std::vector<size_t> Tester<T>::GetOffsets() { +template <typename T, typename U> +const std::vector<size_t> Tester<T,U>::GetOffsets() const { if (full_test_) { return {0, 10}; } else { return {0}; } } // ================================================================================================= -template <> bool Tester<float>::PrecisionSupported() const { return true; } -template <> bool Tester<float2>::PrecisionSupported() const { return true; } -template <> bool Tester<double>::PrecisionSupported() const { - auto extensions = device_.Extensions(); - return (extensions.find(kKhronosDoublePrecision) == std::string::npos) ? false : true; -} -template <> bool Tester<double2>::PrecisionSupported() const { - auto extensions = device_.Extensions(); - return (extensions.find(kKhronosDoublePrecision) == std::string::npos) ? false : true; -} - -// ================================================================================================= - // A test can either pass, be skipped, or fail -template <typename T> -void Tester<T>::ReportPass() { +template <typename T, typename U> +void Tester<T,U>::ReportPass() { num_passed_++; } -template <typename T> -void Tester<T>::ReportSkipped() { +template <typename T, typename U> +void Tester<T,U>::ReportSkipped() { num_skipped_++; } -template <typename T> -void Tester<T>::ReportError(const ErrorLogEntry &error_log_entry) { +template <typename T, typename U> +void Tester<T,U>::ReportError(const ErrorLogEntry &error_log_entry) { error_log_.push_back(error_log_entry); num_failed_++; } @@ -324,8 +253,8 @@ void Tester<T>::ReportError(const ErrorLogEntry &error_log_entry) { // Prints the test-result symbol to screen. This function limits the maximum number of symbols per // line by printing newlines once every so many calls. -template <typename T> -void Tester<T>::PrintTestResult(const std::string &message) { +template <typename T, typename U> +void Tester<T,U>::PrintTestResult(const std::string &message) { if (print_count_ == kResultsPerLine) { print_count_ = 0; fprintf(stdout, "\n "); @@ -336,12 +265,97 @@ void Tester<T>::PrintTestResult(const std::string &message) { } // ================================================================================================= +// Below are the non-member functions (separated because of otherwise required partial class +// template specialization) +// ================================================================================================= + +// Compares two floating point values and returns whether they are within an acceptable error +// margin. This replaces GTest's EXPECT_NEAR(). +template <typename T> +bool TestSimilarity(const T val1, const T val2) { + const auto difference = std::fabs(val1 - val2); + + // Set the allowed error margin for floating-point comparisons + constexpr auto kErrorMarginRelative = 1.0e-2; + constexpr auto kErrorMarginAbsolute = 1.0e-10; + + // Shortcut, handles infinities + if (val1 == val2) { + return true; + } + // The values are zero or very small: the relative error is less meaningful + else if (val1 == 0 || val2 == 0 || difference < static_cast<T>(kErrorMarginAbsolute)) { + return (difference < static_cast<T>(kErrorMarginAbsolute)); + } + // Use relative error + else { + const auto absolute_sum = std::fabs(val1) + std::fabs(val2); + return (difference / absolute_sum) < static_cast<T>(kErrorMarginRelative); + } +} + +// Compiles the default case for non-complex data-types +template bool TestSimilarity<float>(const float, const float); +template bool TestSimilarity<double>(const double, const double); + +// Specialisations for complex data-types +template <> +bool TestSimilarity(const float2 val1, const float2 val2) { + auto real = TestSimilarity(val1.real(), val2.real()); + auto imag = TestSimilarity(val1.imag(), val2.imag()); + return (real && imag); +} +template <> +bool TestSimilarity(const double2 val1, const double2 val2) { + auto real = TestSimilarity(val1.real(), val2.real()); + auto imag = TestSimilarity(val1.imag(), val2.imag()); + return (real && imag); +} + +// ================================================================================================= + +// Retrieves a list of example scalar values, used for the alpha and beta arguments for the various +// routines. This function is specialised for the different data-types. +template <> const std::vector<float> GetExampleScalars(const bool full_test) { + if (full_test) { return {0.0f, 1.0f, 3.14f}; } + else { return {3.14f}; } +} +template <> const std::vector<double> GetExampleScalars(const bool full_test) { + if (full_test) { return {0.0, 1.0, 3.14}; } + else { return {3.14}; } +} +template <> const std::vector<float2> GetExampleScalars(const bool full_test) { + if (full_test) { return {{0.0f, 0.0f}, {1.0f, 1.3f}, {2.42f, 3.14f}}; } + else { return {{2.42f, 3.14f}}; } +} +template <> const std::vector<double2> GetExampleScalars(const bool full_test) { + if (full_test) { return {{0.0, 0.0}, {1.0, 1.3}, {2.42, 3.14}}; } + else { return {{2.42, 3.14}}; } +} + +// ================================================================================================= + +// Returns false is this precision is not supported by the device +template <> bool PrecisionSupported<float>(const Device &) { return true; } +template <> bool PrecisionSupported<float2>(const Device &) { return true; } +template <> bool PrecisionSupported<double>(const Device &device) { + auto extensions = device.Extensions(); + return (extensions.find(kKhronosDoublePrecision) == std::string::npos) ? false : true; +} +template <> bool PrecisionSupported<double2>(const Device &device) { + auto extensions = device.Extensions(); + return (extensions.find(kKhronosDoublePrecision) == std::string::npos) ? false : true; +} + +// ================================================================================================= // Compiles the templated class -template class Tester<float>; -template class Tester<double>; -template class Tester<float2>; -template class Tester<double2>; +template class Tester<float, float>; +template class Tester<double, double>; +template class Tester<float2, float2>; +template class Tester<double2, double2>; +template class Tester<float2, float>; +template class Tester<double2, double>; // ================================================================================================= } // namespace clblast diff --git a/test/correctness/tester.h b/test/correctness/tester.h index 9c4a9e86..93515138 100644 --- a/test/correctness/tester.h +++ b/test/correctness/tester.h @@ -10,6 +10,8 @@ // This file implements the Tester class, providing a test-framework. GTest was used before, but // was not able to handle certain cases (e.g. template type + parameters). This is its (basic) // custom replacement. +// Typename T: the data-type of the routine's memory buffers (==precision) +// Typename U: the data-type of the alpha and beta arguments // // ================================================================================================= @@ -30,7 +32,7 @@ namespace clblast { // ================================================================================================= // See comment at top of file for a description of the class -template <typename T> +template <typename T, typename U> class Tester { public: @@ -43,10 +45,6 @@ class Tester { // Error percentage is not applicable: error was caused by an incorrect status static constexpr auto kStatusError = -1.0f; - // Set the allowed error margin for floating-point comparisons - static constexpr auto kErrorMarginRelative = 1.0e-2; - static constexpr auto kErrorMarginAbsolute = 1.0e-10; - // Constants holding start and end strings for terminal-output in colour const std::string kPrintError{"\x1b[31m"}; const std::string kPrintSuccess{"\x1b[32m"}; @@ -67,7 +65,7 @@ class Tester { StatusCode status_expect; StatusCode status_found; float error_percentage; - Arguments<T> args; + Arguments<U> args; }; // Creates an instance of the tester, running on a particular OpenCL platform and device. It @@ -80,25 +78,13 @@ class Tester { void TestStart(const std::string &test_name, const std::string &test_configuration); void TestEnd(); - // Compares two floating point values for similarity. Allows for a certain relative error margin. - static bool TestSimilarity(const T val1, const T val2); - // Tests either an error count (should be zero) or two error codes (must match) - void TestErrorCount(const size_t errors, const size_t size, const Arguments<T> &args); + void TestErrorCount(const size_t errors, const size_t size, const Arguments<U> &args); void TestErrorCodes(const StatusCode clblas_status, const StatusCode clblast_status, - const Arguments<T> &args); + const Arguments<U> &args); protected: - // Retrieves a list of example scalars of the right type - const std::vector<T> GetExampleScalars(); - - // Retrieves a list of offset values to test - const std::vector<size_t> GetOffsets(); - - // Returns false is this precision is not supported by the device - bool PrecisionSupported() const; - // The help-message std::string help_; @@ -108,6 +94,12 @@ class Tester { Context context_; CommandQueue queue_; + // Whether or not to run the full test-suite or just a smoke test + bool full_test_; + + // Retrieves the offset values to test with + const std::vector<size_t> GetOffsets() const; + private: // Internal methods to report a passed, skipped, or failed test @@ -118,9 +110,6 @@ class Tester { // Prints the error or success symbol to screen void PrintTestResult(const std::string &message); - // Whether or not to run the full test-suite or just a smoke test - bool full_test_; - // Logging and counting occurrences of errors std::vector<ErrorLogEntry> error_log_; size_t num_passed_; @@ -140,6 +129,25 @@ class Tester { }; // ================================================================================================= +// Below are the non-member functions (separated because of otherwise required partial class +// template specialization) +// ================================================================================================= + +// Compares two floating point values and returns whether they are within an acceptable error +// margin. This replaces GTest's EXPECT_NEAR(). +template <typename T> +bool TestSimilarity(const T val1, const T val2); + +// Retrieves a list of example scalar values, used for the alpha and beta arguments for the various +// routines. This function is specialised for the different data-types. +template <typename T> +const std::vector<T> GetExampleScalars(const bool full_test); + +// Returns false is this precision is not supported by the device +template <typename T> +bool PrecisionSupported(const Device &device); + +// ================================================================================================= } // namespace clblast // CLBLAST_TEST_CORRECTNESS_TESTER_H_ diff --git a/test/performance/client.cc b/test/performance/client.cc index 71471dde..676e88e4 100644 --- a/test/performance/client.cc +++ b/test/performance/client.cc @@ -22,10 +22,10 @@ namespace clblast { // ================================================================================================= // Constructor -template <typename T> -Client<T>::Client(const Routine run_routine, const Routine run_reference, - const std::vector<std::string> &options, - const GetMetric get_flops, const GetMetric get_bytes): +template <typename T, typename U> +Client<T,U>::Client(const Routine run_routine, const Routine run_reference, + const std::vector<std::string> &options, + const GetMetric get_flops, const GetMetric get_bytes): run_routine_(run_routine), run_reference_(run_reference), options_(options), @@ -38,10 +38,10 @@ Client<T>::Client(const Routine run_routine, const Routine run_reference, // Parses all arguments available for the CLBlast client testers. Some arguments might not be // applicable, but are searched for anyway to be able to create one common argument parser. All // arguments have a default value in case they are not found. -template <typename T> -Arguments<T> Client<T>::ParseArguments(int argc, char *argv[], const GetMetric default_a_ld, - const GetMetric default_b_ld, const GetMetric default_c_ld) { - auto args = Arguments<T>{}; +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"}; // These are the options which are not for every client: they are optional @@ -58,6 +58,7 @@ Arguments<T> Client<T>::ParseArguments(int argc, char *argv[], const GetMetric d if (o == kArgBTransp) { args.b_transpose = GetArgument(argc, argv, help, kArgBTransp, Transpose::kNo); } if (o == kArgSide) { args.side = GetArgument(argc, argv, help, kArgSide, Side::kLeft); } if (o == kArgTriangle) { args.triangle = GetArgument(argc, argv, help, kArgTriangle, Triangle::kUpper); } + if (o == kArgDiagonal) { args.diagonal = GetArgument(argc, argv, help, kArgDiagonal, Diagonal::kUnit); } // Vector arguments if (o == kArgXInc) { args.x_inc = GetArgument(argc, argv, help, kArgXInc, size_t{1}); } @@ -74,8 +75,8 @@ Arguments<T> Client<T>::ParseArguments(int argc, char *argv[], const GetMetric d if (o == kArgCOffset) { args.c_offset = GetArgument(argc, argv, help, kArgCOffset, size_t{0}); } // Scalar values - if (o == kArgAlpha) { args.alpha = GetArgument(argc, argv, help, kArgAlpha, GetScalar<T>()); } - if (o == kArgBeta) { args.beta = GetArgument(argc, argv, help, kArgBeta, GetScalar<T>()); } + if (o == kArgAlpha) { args.alpha = GetArgument(argc, argv, help, kArgAlpha, GetScalar<U>()); } + if (o == kArgBeta) { args.beta = GetArgument(argc, argv, help, kArgBeta, GetScalar<U>()); } } // These are the options common to all routines @@ -101,8 +102,8 @@ Arguments<T> Client<T>::ParseArguments(int argc, char *argv[], const GetMetric d // ================================================================================================= // This is main performance tester -template <typename T> -void Client<T>::PerformanceTest(Arguments<T> &args, const SetMetric set_sizes) { +template <typename T, typename U> +void Client<T,U>::PerformanceTest(Arguments<U> &args, const SetMetric set_sizes) { // Prints the header of the output table PrintTableHeader(args.silent, options_); @@ -173,10 +174,10 @@ void Client<T>::PerformanceTest(Arguments<T> &args, const SetMetric set_sizes) { // Creates a vector of timing results, filled with execution times of the 'main computation'. The // timing is performed using the milliseconds chrono functions. The function returns the minimum // value found in the vector of timing results. The return value is in milliseconds. -template <typename T> -double Client<T>::TimedExecution(const size_t num_runs, const Arguments<T> &args, - const Buffers &buffers, CommandQueue &queue, - Routine run_blas, const std::string &library_name) { +template <typename T, typename U> +double Client<T,U>::TimedExecution(const size_t num_runs, const Arguments<U> &args, + const Buffers &buffers, CommandQueue &queue, + Routine run_blas, const std::string &library_name) { auto timings = std::vector<double>(num_runs); for (auto &timing: timings) { auto start_time = std::chrono::steady_clock::now(); @@ -197,8 +198,8 @@ double Client<T>::TimedExecution(const size_t num_runs, const Arguments<T> &args // ================================================================================================= // Prints the header of the performance table -template <typename T> -void Client<T>::PrintTableHeader(const bool silent, const std::vector<std::string> &args) { +template <typename T, typename U> +void Client<T,U>::PrintTableHeader(const bool silent, const std::vector<std::string> &args) { if (!silent) { for (auto i=size_t{0}; i<args.size(); ++i) { fprintf(stdout, "%9s ", ""); } fprintf(stdout, " | <-- CLBlast --> | <-- clBLAS --> |\n"); @@ -209,9 +210,9 @@ void Client<T>::PrintTableHeader(const bool silent, const std::vector<std::strin } // Print a performance-result row -template <typename T> -void Client<T>::PrintTableRow(const Arguments<T>& args, const double ms_clblast, - const double ms_clblas) { +template <typename T, typename U> +void Client<T,U>::PrintTableRow(const Arguments<U>& args, const double ms_clblast, + const double ms_clblas) { // Creates a vector of relevant variables auto integers = std::vector<size_t>{}; @@ -224,6 +225,7 @@ void Client<T>::PrintTableRow(const Arguments<T>& args, const double ms_clblast, else if (o == kArgTriangle) { integers.push_back(static_cast<size_t>(args.triangle)); } else if (o == kArgATransp) { integers.push_back(static_cast<size_t>(args.a_transpose)); } else if (o == kArgBTransp) { integers.push_back(static_cast<size_t>(args.b_transpose)); } + else if (o == kArgDiagonal) { integers.push_back(static_cast<size_t>(args.diagonal)); } else if (o == kArgXInc) { integers.push_back(args.x_inc); } else if (o == kArgYInc) { integers.push_back(args.y_inc); } else if (o == kArgXOffset) { integers.push_back(args.x_offset); } @@ -274,10 +276,12 @@ void Client<T>::PrintTableRow(const Arguments<T>& args, const double ms_clblast, // ================================================================================================= // Compiles the templated class -template class Client<float>; -template class Client<double>; -template class Client<float2>; -template class Client<double2>; +template class Client<float,float>; +template class Client<double,double>; +template class Client<float2,float2>; +template class Client<double2,double2>; +template class Client<float2,float>; +template class Client<double2,double>; // ================================================================================================= } // namespace clblast diff --git a/test/performance/client.h b/test/performance/client.h index f9f219d0..c9095967 100644 --- a/test/performance/client.h +++ b/test/performance/client.h @@ -10,6 +10,8 @@ // This class implements the performance-test client. It is generic for all CLBlast routines by // taking a number of routine-specific functions as arguments, such as how to compute buffer sizes // or how to get the FLOPS count. +// Typename T: the data-type of the routine's memory buffers (==precision) +// Typename U: the data-type of the alpha and beta arguments // // This file also provides the common interface to the performance client (see the 'RunClient' // function for details). @@ -32,7 +34,7 @@ namespace clblast { // ================================================================================================= // See comment at top of file for a description of the class -template <typename T> +template <typename T, typename U> class Client { public: @@ -40,9 +42,9 @@ class Client { const cl_device_type kDeviceType = CL_DEVICE_TYPE_ALL; // Shorthand for the routine-specific functions passed to the tester - using Routine = std::function<StatusCode(const Arguments<T>&, const Buffers&, CommandQueue&)>; - using SetMetric = std::function<void(Arguments<T>&)>; - using GetMetric = std::function<size_t(const Arguments<T>&)>; + using Routine = std::function<StatusCode(const Arguments<U>&, const Buffers&, CommandQueue&)>; + using SetMetric = std::function<void(Arguments<U>&)>; + using GetMetric = std::function<size_t(const Arguments<U>&)>; // The constructor Client(const Routine run_routine, const Routine run_reference, @@ -51,24 +53,24 @@ class Client { // Parses all command-line arguments, filling in the arguments structure. If no command-line // argument is given for a particular argument, it is filled in with a default value. - Arguments<T> ParseArguments(int argc, char *argv[], const GetMetric default_a_ld, + Arguments<U> ParseArguments(int argc, char *argv[], const GetMetric default_a_ld, const GetMetric default_b_ld, const GetMetric default_c_ld); // The main client function, setting-up arguments, matrices, OpenCL buffers, etc. After set-up, it // calls the client routines. - void PerformanceTest(Arguments<T> &args, const SetMetric set_sizes); + void PerformanceTest(Arguments<U> &args, const SetMetric set_sizes); private: // Runs a function a given number of times and returns the execution time of the shortest instance - double TimedExecution(const size_t num_runs, const Arguments<T> &args, const Buffers &buffers, + double TimedExecution(const size_t num_runs, const Arguments<U> &args, const Buffers &buffers, CommandQueue &queue, Routine run_blas, const std::string &library_name); // Prints the header of a performance-data table void PrintTableHeader(const bool silent, const std::vector<std::string> &args); // Prints a row of performance data, including results of two libraries - void PrintTableRow(const Arguments<T>& args, const double ms_clblast, const double ms_clblas); + void PrintTableRow(const Arguments<U>& args, const double ms_clblast, const double ms_clblas); // The routine-specific functions passed to the tester const Routine run_routine_; @@ -82,12 +84,12 @@ class Client { // The interface to the performance client. This is a separate function in the header such that it // is automatically compiled for each routine, templated by the parameter "C". -template <typename C, typename T> +template <typename C, typename T, typename U> void RunClient(int argc, char *argv[]) { // Creates a new client - auto client = Client<T>(C::RunRoutine, C::RunReference, C::GetOptions(), - C::GetFlops, C::GetBytes); + auto client = Client<T,U>(C::RunRoutine, C::RunReference, C::GetOptions(), + C::GetFlops, C::GetBytes); // Simple command line argument parser with defaults auto args = client.ParseArguments(argc, argv, C::DefaultLDA, C::DefaultLDB, C::DefaultLDC); diff --git a/test/performance/graphs/xsymm.r b/test/performance/graphs/xsymm.r index 6493f52a..f4b98b30 100644 --- a/test/performance/graphs/xsymm.r +++ b/test/performance/graphs/xsymm.r @@ -19,7 +19,7 @@ source(file.path(dirname(thisfile), "common.r")) # Settings routine_name <- "xsymm" -parameters <- c("-m","-n","-layout","-triangle","-side", +parameters <- c("-m","-n","-layout","-side","-triangle", "-num_steps","-step","-runs","-precision") precision <- 32 @@ -29,7 +29,7 @@ test_names <- list( "multiples of 128 (+1)", "around m=n=512", "around m=n=2048", - "layouts and triangle/side (m=n=1024)", + "layouts and side/triangle (m=n=1024)", "powers of 2" ) @@ -70,7 +70,7 @@ test_xlabels <- list( "matrix sizes (m=n)", "matrix sizes (m=n)", "matrix sizes (m=n)", - "layout (row/col), triangle (up/lo), side (l/r)", + "layout (row/col), side (l/r), triangle (up/lo)", "matrix sizes (m=n)" ) @@ -80,8 +80,8 @@ test_xaxis <- list( c("m", ""), c("m", ""), c("m", ""), - list(1:8, c("row,up,l", "row,up,r", "row,lo,l", "row,lo,r", - "col,up,l", "col,up,r", "col,lo,l", "col,lo,r")), + list(1:8, c("row,l,up", "row,r,up", "row,l,lo", "row,r,lo", + "col,l,up", "col,r,up", "col,l,lo", "col,r,lo")), c("m", "x") ) diff --git a/test/performance/graphs/xtrmm.r b/test/performance/graphs/xtrmm.r new file mode 100644 index 00000000..4a3c2440 --- /dev/null +++ b/test/performance/graphs/xtrmm.r @@ -0,0 +1,127 @@ + +# ================================================================================================== +# This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +# project 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 performance script for the Xtrmm routine +# +# ================================================================================================== + +# Includes the common functions +args <- commandArgs(trailingOnly = FALSE) +thisfile <- (normalizePath(sub("--file=", "", args[grep("--file=", args)]))) +source(file.path(dirname(thisfile), "common.r")) + +# ================================================================================================== + +# Settings +routine_name <- "xtrmm" +parameters <- c("-m","-n","-layout","-side","-triangle","-transA","-diagonal", + "-num_steps","-step","-runs","-precision") +precision <- 32 + +# Sets the names of the test-cases +test_names <- list( + "multiples of 128", + "multiples of 128 (+1)", + "around m=n=512", + "around m=n=2048", + "layouts and side/triangle (m=n=1024)", + "powers of 2" +) + +# Defines the test-cases +test_values <- list( + list(c(128, 128, 0, 0, 0, 0, 0, 16, 128, num_runs, precision)), + list(c(129, 129, 0, 0, 0, 0, 0, 16, 128, num_runs, precision)), + list(c(512, 512, 0, 0, 0, 0, 0, 16, 1, num_runs, precision)), + list(c(2048, 2048, 0, 0, 0, 0, 0, 16, 1, num_runs, precision)), + list( + c(1024, 1024, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 0, 0, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 0, 1, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 0, 1, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 1, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 1, 0, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 1, 1, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 1, 1, 1, 1, 0, num_runs, precision), + + c(1024, 1024, 0, 1, 0, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 1, 0, 0, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 1, 0, 1, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 1, 0, 1, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 1, 1, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 1, 1, 0, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 1, 1, 1, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 1, 1, 1, 1, 1, 0, num_runs, precision), + + c(1024, 1024, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 0, 0, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 0, 1, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 0, 1, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 1, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 1, 0, 1, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 1, 1, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 1, 1, 1, 1, 0, num_runs, precision), + + c(1024, 1024, 1, 1, 0, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 1, 1, 0, 0, 1, 1, 0, num_runs, precision), + c(1024, 1024, 1, 1, 0, 1, 0, 1, 0, num_runs, precision), + c(1024, 1024, 1, 1, 0, 1, 1, 1, 0, num_runs, precision), + c(1024, 1024, 1, 1, 1, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 1, 1, 1, 0, 1, 1, 0, num_runs, precision), + c(1024, 1024, 1, 1, 1, 1, 0, 1, 0, num_runs, precision), + c(1024, 1024, 1, 1, 1, 1, 1, 1, 0, num_runs, precision) + ), + list( + c(8, 8, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(16, 16, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(32, 32, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(64, 64, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(128, 128, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(256, 256, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(512, 512, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(1024, 1024, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(2048, 2048, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(4096, 4096, 0, 0, 0, 0, 0, 1, 0, num_runs, precision), + c(8192, 8192, 0, 0, 0, 0, 0, 1, 0, num_runs, precision) + ) +) + +# Defines the x-labels corresponding to the test-cases +test_xlabels <- list( + "matrix sizes (m=n)", + "matrix sizes (m=n)", + "matrix sizes (m=n)", + "matrix sizes (m=n)", + "layout (row/col), side (l/r), triangle (up/lo), transA (n/y), diag (u/nu)", + "matrix sizes (m=n)" +) + +# Defines the x-axis of the test-cases +test_xaxis <- list( + c("m", ""), + c("m", ""), + c("m", ""), + c("m", ""), + list(1:32, c("row,l,up,n,u", "row,l,up,n,nu", "row,l,up,y,u", "row,l,up,y,nu", + "row,r,up,n,u", "row,r,up,n,nu", "row,r,up,y,u", "row,r,up,y,nu", + "row,l,lo,n,u", "row,l,lo,n,nu", "row,l,lo,y,u", "row,l,lo,y,nu", + "row,r,lo,n,u", "row,r,lo,n,nu", "row,r,lo,y,u", "row,r,lo,y,nu", + "col,l,up,n,u", "col,l,up,n,nu", "col,l,up,y,u", "col,l,up,y,nu", + "col,r,up,n,u", "col,r,up,n,nu", "col,r,up,y,u", "col,r,up,y,nu", + "col,l,lo,n,u", "col,l,lo,n,nu", "col,l,lo,y,u", "col,l,lo,y,nu", + "col,r,lo,n,u", "col,r,lo,n,nu", "col,r,lo,y,u", "col,r,lo,y,nu")), + c("m", "x") +) + +# ================================================================================================== + +# Start the script +main(routine_name=routine_name, precision=precision, test_names=test_names, test_values=test_values, + test_xlabels=test_xlabels, test_xaxis=test_xaxis, metric_gflops=TRUE) + +# ==================================================================================================
\ No newline at end of file diff --git a/test/performance/routines/xaxpy.cc b/test/performance/routines/xaxpy.cc index 3ced80ed..6a2b96c1 100644 --- a/test/performance/routines/xaxpy.cc +++ b/test/performance/routines/xaxpy.cc @@ -16,19 +16,23 @@ // ================================================================================================= +// 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::TestXaxpy<float>, float>(argc, argv); break; + clblast::RunClient<clblast::TestXaxpy<float>, float, float>(argc, argv); break; case clblast::Precision::kDouble: - clblast::RunClient<clblast::TestXaxpy<double>, double>(argc, argv); break; + clblast::RunClient<clblast::TestXaxpy<double>, double, double>(argc, argv); break; case clblast::Precision::kComplexSingle: - clblast::RunClient<clblast::TestXaxpy<clblast::float2>, clblast::float2>(argc, argv); break; + clblast::RunClient<clblast::TestXaxpy<float2>, float2, float2>(argc, argv); break; case clblast::Precision::kComplexDouble: - clblast::RunClient<clblast::TestXaxpy<clblast::double2>, clblast::double2>(argc, argv); break; + clblast::RunClient<clblast::TestXaxpy<double2>, double2, double2>(argc, argv); break; } return 0; } diff --git a/test/performance/routines/xgemm.cc b/test/performance/routines/xgemm.cc index 36c74b9a..9a02e595 100644 --- a/test/performance/routines/xgemm.cc +++ b/test/performance/routines/xgemm.cc @@ -16,19 +16,23 @@ // ================================================================================================= +// 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::TestXgemm<float>, float>(argc, argv); break; + clblast::RunClient<clblast::TestXgemm<float>, float, float>(argc, argv); break; case clblast::Precision::kDouble: - clblast::RunClient<clblast::TestXgemm<double>, double>(argc, argv); break; + clblast::RunClient<clblast::TestXgemm<double>, double, double>(argc, argv); break; case clblast::Precision::kComplexSingle: - clblast::RunClient<clblast::TestXgemm<clblast::float2>, clblast::float2>(argc, argv); break; + clblast::RunClient<clblast::TestXgemm<float2>, float2, float2>(argc, argv); break; case clblast::Precision::kComplexDouble: - clblast::RunClient<clblast::TestXgemm<clblast::double2>, clblast::double2>(argc, argv); break; + clblast::RunClient<clblast::TestXgemm<double2>, double2, double2>(argc, argv); break; } return 0; } diff --git a/test/performance/routines/xgemv.cc b/test/performance/routines/xgemv.cc index 183dd4a1..6f69ef21 100644 --- a/test/performance/routines/xgemv.cc +++ b/test/performance/routines/xgemv.cc @@ -16,19 +16,23 @@ // ================================================================================================= +// 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::TestXgemv<float>, float>(argc, argv); break; + clblast::RunClient<clblast::TestXgemv<float>, float, float>(argc, argv); break; case clblast::Precision::kDouble: - clblast::RunClient<clblast::TestXgemv<double>, double>(argc, argv); break; + clblast::RunClient<clblast::TestXgemv<double>, double, double>(argc, argv); break; case clblast::Precision::kComplexSingle: - clblast::RunClient<clblast::TestXgemv<clblast::float2>, clblast::float2>(argc, argv); break; + clblast::RunClient<clblast::TestXgemv<float2>, float2, float2>(argc, argv); break; case clblast::Precision::kComplexDouble: - clblast::RunClient<clblast::TestXgemv<clblast::double2>, clblast::double2>(argc, argv); break; + clblast::RunClient<clblast::TestXgemv<double2>, double2, double2>(argc, argv); break; } return 0; } diff --git a/test/performance/routines/xhemm.cc b/test/performance/routines/xhemm.cc new file mode 100644 index 00000000..34798d8d --- /dev/null +++ b/test/performance/routines/xhemm.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 Xhemm command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/xhemm.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::TestXhemm<float2>, float2, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXhemm<double2>, double2, double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/test/performance/routines/xher2k.cc b/test/performance/routines/xher2k.cc new file mode 100644 index 00000000..1b505737 --- /dev/null +++ b/test/performance/routines/xher2k.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 Xher2k command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/xher2k.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::TestXher2k<float2,float>, float2, float>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXher2k<double2,double>, double2, double>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/test/performance/routines/xherk.cc b/test/performance/routines/xherk.cc new file mode 100644 index 00000000..ce18152e --- /dev/null +++ b/test/performance/routines/xherk.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 Xherk command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/xherk.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::TestXherk<float2,float>, float2, float>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXherk<double2,double>, double2, double>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/test/performance/routines/xsymm.cc b/test/performance/routines/xsymm.cc index 0c7f5e1e..8738ceda 100644 --- a/test/performance/routines/xsymm.cc +++ b/test/performance/routines/xsymm.cc @@ -16,19 +16,23 @@ // ================================================================================================= +// 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::TestXsymm<float>, float>(argc, argv); break; + clblast::RunClient<clblast::TestXsymm<float>, float, float>(argc, argv); break; case clblast::Precision::kDouble: - clblast::RunClient<clblast::TestXsymm<double>, double>(argc, argv); break; + clblast::RunClient<clblast::TestXsymm<double>, double, double>(argc, argv); break; case clblast::Precision::kComplexSingle: - clblast::RunClient<clblast::TestXsymm<clblast::float2>, clblast::float2>(argc, argv); break; + clblast::RunClient<clblast::TestXsymm<float2>, float2, float2>(argc, argv); break; case clblast::Precision::kComplexDouble: - clblast::RunClient<clblast::TestXsymm<clblast::double2>, clblast::double2>(argc, argv); break; + clblast::RunClient<clblast::TestXsymm<double2>, double2, double2>(argc, argv); break; } return 0; } diff --git a/test/performance/routines/xsyr2k.cc b/test/performance/routines/xsyr2k.cc index 63b50df6..e4c76229 100644 --- a/test/performance/routines/xsyr2k.cc +++ b/test/performance/routines/xsyr2k.cc @@ -16,19 +16,23 @@ // ================================================================================================= +// 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::TestXsyr2k<float>, float>(argc, argv); break; + clblast::RunClient<clblast::TestXsyr2k<float>, float, float>(argc, argv); break; case clblast::Precision::kDouble: - clblast::RunClient<clblast::TestXsyr2k<double>, double>(argc, argv); break; + clblast::RunClient<clblast::TestXsyr2k<double>, double, double>(argc, argv); break; case clblast::Precision::kComplexSingle: - clblast::RunClient<clblast::TestXsyr2k<clblast::float2>, clblast::float2>(argc, argv); break; + clblast::RunClient<clblast::TestXsyr2k<float2>, float2, float2>(argc, argv); break; case clblast::Precision::kComplexDouble: - clblast::RunClient<clblast::TestXsyr2k<clblast::double2>, clblast::double2>(argc, argv); break; + clblast::RunClient<clblast::TestXsyr2k<double2>, double2, double2>(argc, argv); break; } return 0; } diff --git a/test/performance/routines/xsyrk.cc b/test/performance/routines/xsyrk.cc index 9022d4f8..53fecb69 100644 --- a/test/performance/routines/xsyrk.cc +++ b/test/performance/routines/xsyrk.cc @@ -16,19 +16,23 @@ // ================================================================================================= +// 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::TestXsyrk<float>, float>(argc, argv); break; + clblast::RunClient<clblast::TestXsyrk<float>, float, float>(argc, argv); break; case clblast::Precision::kDouble: - clblast::RunClient<clblast::TestXsyrk<double>, double>(argc, argv); break; + clblast::RunClient<clblast::TestXsyrk<double>, double, double>(argc, argv); break; case clblast::Precision::kComplexSingle: - clblast::RunClient<clblast::TestXsyrk<clblast::float2>, clblast::float2>(argc, argv); break; + clblast::RunClient<clblast::TestXsyrk<float2>, float2, float2>(argc, argv); break; case clblast::Precision::kComplexDouble: - clblast::RunClient<clblast::TestXsyrk<clblast::double2>, clblast::double2>(argc, argv); break; + clblast::RunClient<clblast::TestXsyrk<double2>, double2, double2>(argc, argv); break; } return 0; } diff --git a/test/performance/routines/xtrmm.cc b/test/performance/routines/xtrmm.cc new file mode 100644 index 00000000..2ab9ce77 --- /dev/null +++ b/test/performance/routines/xtrmm.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 Xtrmm command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/xtrmm.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::TestXtrmm<float>, float, float>(argc, argv); break; + case clblast::Precision::kDouble: + clblast::RunClient<clblast::TestXtrmm<double>, double, double>(argc, argv); break; + case clblast::Precision::kComplexSingle: + clblast::RunClient<clblast::TestXtrmm<float2>, float2, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient<clblast::TestXtrmm<double2>, double2, double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= diff --git a/test/routines/xhemm.h b/test/routines/xhemm.h new file mode 100644 index 00000000..75878b06 --- /dev/null +++ b/test/routines/xhemm.h @@ -0,0 +1,134 @@ + +// ================================================================================================= +// 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 Xhemm 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_XHEMM_H_ +#define CLBLAST_TEST_ROUTINES_XHEMM_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 TestXhemm { + public: + + // The list of arguments relevant for this routine + static std::vector<std::string> GetOptions() { + return {kArgM, kArgN, + kArgLayout, kArgSide, kArgTriangle, + kArgALeadDim, kArgBLeadDim, kArgCLeadDim, + kArgAOffset, kArgBOffset, kArgCOffset, + kArgAlpha, kArgBeta}; + } + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeA(const Arguments<T> &args) { + size_t k_value = (args.side == Side::kLeft) ? args.m : args.n; + auto a_rotated = (args.layout == Layout::kRowMajor); + auto a_two = (a_rotated) ? args.m : k_value; + return a_two * args.a_ld + args.a_offset; + } + static size_t GetSizeB(const Arguments<T> &args) { + size_t k_value = (args.side == Side::kLeft) ? args.m : args.n; + auto b_rotated = (args.layout == Layout::kRowMajor); + auto b_two = (b_rotated) ? k_value : args.n; + return b_two * args.b_ld + args.b_offset; + } + static size_t GetSizeC(const Arguments<T> &args) { + auto c_rotated = (args.layout == Layout::kRowMajor); + auto c_two = (c_rotated) ? args.m : args.n; + return c_two * args.c_ld + args.c_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments<T> &args) { + args.a_size = GetSizeA(args); + args.b_size = GetSizeB(args); + args.c_size = GetSizeC(args); + } + + // Describes what the default values of the leading dimensions of the matrices are + static size_t DefaultLDA(const Arguments<T> &args) { return args.m; } + static size_t DefaultLDB(const Arguments<T> &args) { return args.n; } + static size_t DefaultLDC(const Arguments<T> &args) { return args.n; } + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments<T> &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = Hemm(args.layout, args.side, args.triangle, + args.m, args.n, args.alpha, + buffers.a_mat(), args.a_offset, args.a_ld, + buffers.b_mat(), args.b_offset, args.b_ld, args.beta, + buffers.c_mat(), args.c_offset, args.c_ld, + &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 &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXhemm(static_cast<clblasOrder>(args.layout), + static_cast<clblasSide>(args.side), + static_cast<clblasUplo>(args.triangle), + args.m, args.n, args.alpha, + buffers.a_mat(), args.a_offset, args.a_ld, + buffers.b_mat(), args.b_offset, args.b_ld, args.beta, + buffers.c_mat(), args.c_offset, args.c_ld, + 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 &buffers, + CommandQueue &queue) { + std::vector<T> result(args.c_size, static_cast<T>(0)); + buffers.c_mat.ReadBuffer(queue, args.c_size*sizeof(T), result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<T> &args) { return args.m; } + static size_t ResultID2(const Arguments<T> &args) { return args.n; } + static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t id2) { + return (args.layout == Layout::kRowMajor) ? + id1*args.c_ld + id2 + args.c_offset: + id2*args.c_ld + id1 + args.c_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<T> &args) { + return 2 * args.m * args.n * args.m; + } + static size_t GetBytes(const Arguments<T> &args) { + return (args.m*args.m + args.m*args.n + 2*args.m*args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XHEMM_H_ +#endif diff --git a/test/routines/xher2k.h b/test/routines/xher2k.h new file mode 100644 index 00000000..f13e8a62 --- /dev/null +++ b/test/routines/xher2k.h @@ -0,0 +1,132 @@ + +// ================================================================================================= +// 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 Xher2k 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_XHER2K_H_ +#define CLBLAST_TEST_ROUTINES_XHER2K_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, typename U> +class TestXher2k { + public: + + // The list of arguments relevant for this routine + static std::vector<std::string> GetOptions() { + return {kArgN, kArgK, + kArgLayout, kArgTriangle, kArgATransp, + kArgALeadDim, kArgBLeadDim, kArgCLeadDim, + kArgAOffset, kArgBOffset, kArgCOffset, + kArgAlpha, kArgBeta}; + } + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeA(const Arguments<U> &args) { + auto a_rotated = (args.layout == Layout::kColMajor && args.a_transpose != Transpose::kNo) || + (args.layout == Layout::kRowMajor && args.a_transpose == Transpose::kNo); + auto a_two = (a_rotated) ? args.n : args.k; + return a_two * args.a_ld + args.a_offset; + } + static size_t GetSizeB(const Arguments<U> &args) { + auto b_rotated = (args.layout == Layout::kColMajor && args.b_transpose != Transpose::kNo) || + (args.layout == Layout::kRowMajor && args.b_transpose == Transpose::kNo); + auto b_two = (b_rotated) ? args.n : args.k; + return b_two * args.b_ld + args.b_offset; + } + static size_t GetSizeC(const Arguments<U> &args) { + return args.n * args.c_ld + args.c_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments<U> &args) { + args.a_size = GetSizeA(args); + args.b_size = GetSizeB(args); + args.c_size = GetSizeC(args); + } + + // Describes what the default values of the leading dimensions of the matrices are + static size_t DefaultLDA(const Arguments<U> &args) { return args.k; } + static size_t DefaultLDB(const Arguments<U> &args) { return args.k; } + static size_t DefaultLDC(const Arguments<U> &args) { return args.n; } + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments<U> &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto alpha2 = T{args.alpha, args.alpha}; + auto status = Her2k(args.layout, args.triangle, args.a_transpose, + args.n, args.k, alpha2, + buffers.a_mat(), args.a_offset, args.a_ld, + buffers.b_mat(), args.b_offset, args.b_ld, args.beta, + buffers.c_mat(), args.c_offset, args.c_ld, + &queue_plain, &event); + clWaitForEvents(1, &event); + return status; + } + + // Describes how to run the clBLAS routine (for correctness/performance comparison) + static StatusCode RunReference(const Arguments<U> &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto alpha2 = T{args.alpha, args.alpha}; + auto status = clblasXher2k(static_cast<clblasOrder>(args.layout), + static_cast<clblasUplo>(args.triangle), + static_cast<clblasTranspose>(args.a_transpose), + args.n, args.k, alpha2, + buffers.a_mat(), args.a_offset, args.a_ld, + buffers.b_mat(), args.b_offset, args.b_ld, args.beta, + buffers.c_mat(), args.c_offset, args.c_ld, + 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<U> &args, Buffers &buffers, + CommandQueue &queue) { + std::vector<T> result(args.c_size, static_cast<T>(0)); + buffers.c_mat.ReadBuffer(queue, args.c_size*sizeof(T), result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<U> &args) { return args.n; } + static size_t ResultID2(const Arguments<U> &args) { return args.n; } + static size_t GetResultIndex(const Arguments<U> &args, const size_t id1, const size_t id2) { + return id1*args.c_ld + id2 + args.c_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<U> &args) { + return 2 * args.n * args.n * args.k; + } + static size_t GetBytes(const Arguments<U> &args) { + return (args.n*args.k + args.n*args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XHER2K_H_ +#endif diff --git a/test/routines/xherk.h b/test/routines/xherk.h new file mode 100644 index 00000000..780b9b52 --- /dev/null +++ b/test/routines/xherk.h @@ -0,0 +1,121 @@ + +// ================================================================================================= +// 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 Xherk 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_XHERK_H_ +#define CLBLAST_TEST_ROUTINES_XHERK_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, typename U> +class TestXherk { + public: + + // The list of arguments relevant for this routine + static std::vector<std::string> GetOptions() { + return {kArgN, kArgK, + kArgLayout, kArgTriangle, kArgATransp, + kArgALeadDim, kArgCLeadDim, + kArgAOffset, kArgCOffset, + kArgAlpha, kArgBeta}; + } + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeA(const Arguments<U> &args) { + auto a_rotated = (args.layout == Layout::kColMajor && args.a_transpose != Transpose::kNo) || + (args.layout == Layout::kRowMajor && args.a_transpose == Transpose::kNo); + auto a_two = (a_rotated) ? args.n : args.k; + return a_two * args.a_ld + args.a_offset; + } + static size_t GetSizeC(const Arguments<U> &args) { + return args.n * args.c_ld + args.c_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments<U> &args) { + args.a_size = GetSizeA(args); + args.c_size = GetSizeC(args); + } + + // Describes what the default values of the leading dimensions of the matrices are + static size_t DefaultLDA(const Arguments<U> &args) { return args.k; } + static size_t DefaultLDB(const Arguments<U> &) { return 1; } // N/A for this routine + static size_t DefaultLDC(const Arguments<U> &args) { return args.n; } + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments<U> &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = Herk(args.layout, args.triangle, args.a_transpose, + args.n, args.k, args.alpha, + buffers.a_mat(), args.a_offset, args.a_ld, args.beta, + buffers.c_mat(), args.c_offset, args.c_ld, + &queue_plain, &event); + clWaitForEvents(1, &event); + return status; + } + + // Describes how to run the clBLAS routine (for correctness/performance comparison) + static StatusCode RunReference(const Arguments<U> &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXherk(static_cast<clblasOrder>(args.layout), + static_cast<clblasUplo>(args.triangle), + static_cast<clblasTranspose>(args.a_transpose), + args.n, args.k, args.alpha, + buffers.a_mat(), args.a_offset, args.a_ld, args.beta, + buffers.c_mat(), args.c_offset, args.c_ld, + 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<U> &args, Buffers &buffers, + CommandQueue &queue) { + std::vector<T> result(args.c_size, static_cast<T>(0)); + buffers.c_mat.ReadBuffer(queue, args.c_size*sizeof(T), result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<U> &args) { return args.n; } + static size_t ResultID2(const Arguments<U> &args) { return args.n; } + static size_t GetResultIndex(const Arguments<U> &args, const size_t id1, const size_t id2) { + return id1*args.c_ld + id2 + args.c_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<U> &args) { + return args.n * args.n * args.k; + } + static size_t GetBytes(const Arguments<U> &args) { + return (args.n*args.k + args.n*args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XHERK_H_ +#endif diff --git a/test/routines/xtrmm.h b/test/routines/xtrmm.h new file mode 100644 index 00000000..7b7e7af1 --- /dev/null +++ b/test/routines/xtrmm.h @@ -0,0 +1,127 @@ + +// ================================================================================================= +// 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 Xtrmm 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_XTRMM_H_ +#define CLBLAST_TEST_ROUTINES_XTRMM_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 TestXtrmm { + public: + + // The list of arguments relevant for this routine + static std::vector<std::string> GetOptions() { + return {kArgM, kArgN, + kArgLayout, kArgSide, kArgTriangle, kArgATransp, kArgDiagonal, + kArgALeadDim, kArgBLeadDim, + kArgAOffset, kArgBOffset, + kArgAlpha}; + } + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeA(const Arguments<T> &args) { + auto k = (args.side == Side::kLeft) ? args.m : args.n; + return k * args.a_ld + args.a_offset; + } + static size_t GetSizeB(const Arguments<T> &args) { + auto b_rotated = (args.layout == Layout::kRowMajor); + auto b_two = (b_rotated) ? args.m : args.n; + return b_two * args.b_ld + args.b_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments<T> &args) { + args.a_size = GetSizeA(args); + args.b_size = GetSizeB(args); + } + + // Describes what the default values of the leading dimensions of the matrices are + static size_t DefaultLDA(const Arguments<T> &args) { return args.m; } + static size_t DefaultLDB(const Arguments<T> &args) { return args.n; } + static size_t DefaultLDC(const Arguments<T> &) { return 1; } // N/A for this routine + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments<T> &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = Trmm(args.layout, args.side, args.triangle, args.a_transpose, args.diagonal, + args.m, args.n, args.alpha, + buffers.a_mat(), args.a_offset, args.a_ld, + buffers.b_mat(), args.b_offset, args.b_ld, + &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 &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXtrmm(static_cast<clblasOrder>(args.layout), + static_cast<clblasSide>(args.side), + static_cast<clblasUplo>(args.triangle), + static_cast<clblasTranspose>(args.a_transpose), + static_cast<clblasDiag>(args.diagonal), + args.m, args.n, args.alpha, + buffers.a_mat(), args.a_offset, args.a_ld, + buffers.b_mat(), args.b_offset, args.b_ld, + 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 &buffers, + CommandQueue &queue) { + std::vector<T> result(args.b_size, static_cast<T>(0)); + buffers.b_mat.ReadBuffer(queue, args.b_size*sizeof(T), result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments<T> &args) { return args.m; } + static size_t ResultID2(const Arguments<T> &args) { return args.n; } + static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t id2) { + return (args.layout == Layout::kRowMajor) ? + id1*args.b_ld + id2 + args.b_offset: + id2*args.b_ld + id1 + args.b_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments<T> &args) { + auto k = (args.side == Side::kLeft) ? args.m : args.n; + return args.m * args.n * k; + } + static size_t GetBytes(const Arguments<T> &args) { + auto k = (args.side == Side::kLeft) ? args.m : args.n; + return (k*k + 2*args.m*args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XTRMM_H_ +#endif diff --git a/test/wrapper_clblas.h b/test/wrapper_clblas.h index d106e4c3..4aaf3705 100644 --- a/test/wrapper_clblas.h +++ b/test/wrapper_clblas.h @@ -76,33 +76,33 @@ clblasStatus clblasXaxpy( // Calls {clblasSgemv, clblasDgemv, clblasCgemv, clblasZgemv} with the arguments forwarded. clblasStatus clblasXgemv( - clblasOrder layout, clblasTranspose tran_a, size_t m, size_t n, float alpha, + clblasOrder layout, clblasTranspose a_transpose, size_t m, size_t n, float alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem x_vec, size_t x_offset, size_t x_inc, float beta, const cl_mem y_vec, size_t y_offset, size_t y_inc, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { - return clblasSgemv(layout, tran_a, m, n, alpha, + return clblasSgemv(layout, a_transpose, m, n, alpha, a_mat, a_offset, a_ld, x_vec, x_offset, static_cast<int>(x_inc), beta, y_vec, y_offset, static_cast<int>(y_inc), num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXgemv( - clblasOrder layout, clblasTranspose tran_a, size_t m, size_t n, double alpha, + clblasOrder layout, clblasTranspose a_transpose, size_t m, size_t n, double alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem x_vec, size_t x_offset, size_t x_inc, double beta, const cl_mem y_vec, size_t y_offset, size_t y_inc, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { - return clblasDgemv(layout, tran_a, m, n, alpha, + return clblasDgemv(layout, a_transpose, m, n, alpha, a_mat, a_offset, a_ld, x_vec, x_offset, static_cast<int>(x_inc), beta, y_vec, y_offset, static_cast<int>(y_inc), num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXgemv( - clblasOrder layout, clblasTranspose tran_a, size_t m, size_t n, float2 alpha, + clblasOrder layout, clblasTranspose a_transpose, size_t m, size_t n, float2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem x_vec, size_t x_offset, size_t x_inc, float2 beta, const cl_mem y_vec, size_t y_offset, size_t y_inc, @@ -110,14 +110,14 @@ clblasStatus clblasXgemv( cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto cl_alpha = cl_float2{{alpha.real(), alpha.imag()}}; auto cl_beta = cl_float2{{beta.real(), beta.imag()}}; - return clblasCgemv(layout, tran_a, m, n, cl_alpha, + return clblasCgemv(layout, a_transpose, m, n, cl_alpha, a_mat, a_offset, a_ld, x_vec, x_offset, static_cast<int>(x_inc), cl_beta, y_vec, y_offset, static_cast<int>(y_inc), num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXgemv( - clblasOrder layout, clblasTranspose tran_a, size_t m, size_t n, double2 alpha, + clblasOrder layout, clblasTranspose a_transpose, size_t m, size_t n, double2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem x_vec, size_t x_offset, size_t x_inc, double2 beta, const cl_mem y_vec, size_t y_offset, size_t y_inc, @@ -125,7 +125,7 @@ clblasStatus clblasXgemv( cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto cl_alpha = cl_double2{{alpha.real(), alpha.imag()}}; auto cl_beta = cl_double2{{beta.real(), beta.imag()}}; - return clblasZgemv(layout, tran_a, m, n, cl_alpha, + return clblasZgemv(layout, a_transpose, m, n, cl_alpha, a_mat, a_offset, a_ld, x_vec, x_offset, static_cast<int>(x_inc), cl_beta, y_vec, y_offset, static_cast<int>(y_inc), @@ -137,14 +137,14 @@ clblasStatus clblasXgemv( // This calls {clblasSgemm, clblasDgemm, clblasCgemm, clblasZgemm} with the arguments forwarded. clblasStatus clblasXgemm( - clblasOrder layout, clblasTranspose tran_a, clblasTranspose tran_b, + clblasOrder layout, clblasTranspose a_transpose, clblasTranspose b_transpose, size_t m, size_t n, size_t k, float alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, float beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { - return clblasSgemm(layout, tran_a, tran_b, + return clblasSgemm(layout, a_transpose, b_transpose, m, n, k, alpha, a_mat, a_offset, a_ld, b_mat, b_offset, b_ld, beta, @@ -152,14 +152,14 @@ clblasStatus clblasXgemm( num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXgemm( - clblasOrder layout, clblasTranspose tran_a, clblasTranspose tran_b, + clblasOrder layout, clblasTranspose a_transpose, clblasTranspose b_transpose, size_t m, size_t n, size_t k, double alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, double beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { - return clblasDgemm(layout, tran_a, tran_b, + return clblasDgemm(layout, a_transpose, b_transpose, m, n, k, alpha, a_mat, a_offset, a_ld, b_mat, b_offset, b_ld, beta, @@ -167,16 +167,16 @@ clblasStatus clblasXgemm( num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXgemm( - clblasOrder layout, clblasTranspose tran_a, clblasTranspose tran_b, + clblasOrder layout, clblasTranspose a_transpose, clblasTranspose b_transpose, size_t m, size_t n, size_t k, float2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, float2 beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto cl_alpha = cl_float2{{alpha.real(), alpha.imag()}}; auto cl_beta = cl_float2{{beta.real(), beta.imag()}}; - return clblasCgemm(layout, tran_a, tran_b, + return clblasCgemm(layout, a_transpose, b_transpose, m, n, k, cl_alpha, a_mat, a_offset, a_ld, b_mat, b_offset, b_ld, cl_beta, @@ -184,16 +184,16 @@ clblasStatus clblasXgemm( num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXgemm( - clblasOrder layout, clblasTranspose tran_a, clblasTranspose tran_b, + clblasOrder layout, clblasTranspose a_transpose, clblasTranspose b_transpose, size_t m, size_t n, size_t k, double2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, double2 beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto cl_alpha = cl_double2{{alpha.real(), alpha.imag()}}; auto cl_beta = cl_double2{{beta.real(), beta.imag()}}; - return clblasZgemm(layout, tran_a, tran_b, + return clblasZgemm(layout, a_transpose, b_transpose, m, n, k, cl_alpha, a_mat, a_offset, a_ld, b_mat, b_offset, b_ld, cl_beta, @@ -207,7 +207,7 @@ clblasStatus clblasXsymm( size_t m, size_t n, float alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, float beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { return clblasSsymm(layout, side, triangle, @@ -222,7 +222,7 @@ clblasStatus clblasXsymm( size_t m, size_t n, double alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, double beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { return clblasDsymm(layout, side, triangle, @@ -237,7 +237,7 @@ clblasStatus clblasXsymm( size_t m, size_t n, float2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, float2 beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, 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()}}; @@ -254,7 +254,7 @@ clblasStatus clblasXsymm( size_t m, size_t n, double2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, double2 beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, 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()}}; @@ -267,74 +267,138 @@ clblasStatus clblasXsymm( num_queues, queues, num_wait_events, wait_events, events); } +// This calls {clblasChemm, clblasZhemm} with the arguments forwarded. +clblasStatus clblasXhemm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + size_t m, size_t n, float2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, float2 beta, + const cl_mem c_mat, size_t c_offset, size_t c_ld, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto cl_alpha = cl_float2{{alpha.real(), alpha.imag()}}; + auto cl_beta = cl_float2{{beta.real(), beta.imag()}}; + return clblasChemm(layout, side, triangle, + m, n, cl_alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, cl_beta, + c_mat, c_offset, c_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXhemm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + size_t m, size_t n, double2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, double2 beta, + const cl_mem c_mat, size_t c_offset, size_t c_ld, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto cl_alpha = cl_double2{{alpha.real(), alpha.imag()}}; + auto cl_beta = cl_double2{{beta.real(), beta.imag()}}; + return clblasZhemm(layout, side, triangle, + m, n, cl_alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, cl_beta, + c_mat, c_offset, c_ld, + num_queues, queues, num_wait_events, wait_events, events); +} + // This calls {clblasSsyrk, clblasDsyrk, clblasCsyrk, clblasZsyrk} with the arguments forwarded. clblasStatus clblasXsyrk( - clblasOrder layout, clblasUplo triangle, clblasTranspose tran_a, + clblasOrder layout, clblasUplo triangle, clblasTranspose a_transpose, size_t n, size_t k, float alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, float beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { - return clblasSsyrk(layout, triangle, tran_a, + return clblasSsyrk(layout, triangle, a_transpose, n, k, alpha, a_mat, a_offset, a_ld, beta, c_mat, c_offset, c_ld, num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXsyrk( - clblasOrder layout, clblasUplo triangle, clblasTranspose tran_a, + clblasOrder layout, clblasUplo triangle, clblasTranspose a_transpose, size_t n, size_t k, double alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, double beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { - return clblasDsyrk(layout, triangle, tran_a, + return clblasDsyrk(layout, triangle, a_transpose, n, k, alpha, a_mat, a_offset, a_ld, beta, c_mat, c_offset, c_ld, num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXsyrk( - clblasOrder layout, clblasUplo triangle, clblasTranspose tran_a, + clblasOrder layout, clblasUplo triangle, clblasTranspose a_transpose, size_t n, size_t k, float2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, float2 beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto cl_alpha = cl_float2{{alpha.real(), alpha.imag()}}; auto cl_beta = cl_float2{{beta.real(), beta.imag()}}; - return clblasCsyrk(layout, triangle, tran_a, + return clblasCsyrk(layout, triangle, a_transpose, n, k, cl_alpha, a_mat, a_offset, a_ld, cl_beta, c_mat, c_offset, c_ld, num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXsyrk( - clblasOrder layout, clblasUplo triangle, clblasTranspose tran_a, + clblasOrder layout, clblasUplo triangle, clblasTranspose a_transpose, size_t n, size_t k, double2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, double2 beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto cl_alpha = cl_double2{{alpha.real(), alpha.imag()}}; auto cl_beta = cl_double2{{beta.real(), beta.imag()}}; - return clblasZsyrk(layout, triangle, tran_a, + return clblasZsyrk(layout, triangle, a_transpose, n, k, cl_alpha, a_mat, a_offset, a_ld, cl_beta, c_mat, c_offset, c_ld, num_queues, queues, num_wait_events, wait_events, events); } +// This calls {clblasCherk, clblasZherk} with the arguments forwarded. +clblasStatus clblasXherk( + clblasOrder layout, clblasUplo triangle, clblasTranspose a_transpose, + size_t n, size_t k, float alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, float beta, + const cl_mem c_mat, size_t c_offset, size_t c_ld, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasCherk(layout, triangle, a_transpose, + n, k, alpha, + a_mat, a_offset, a_ld, beta, + c_mat, c_offset, c_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXherk( + clblasOrder layout, clblasUplo triangle, clblasTranspose a_transpose, + size_t n, size_t k, double alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, double beta, + const cl_mem c_mat, size_t c_offset, size_t c_ld, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasZherk(layout, triangle, a_transpose, + n, k, alpha, + a_mat, a_offset, a_ld, beta, + c_mat, c_offset, c_ld, + num_queues, queues, num_wait_events, wait_events, events); +} + // This calls {clblasSsyr2k, clblasDsyr2k, clblasCsyr2k, clblasZsyr2k} with the arguments forwarded. clblasStatus clblasXsyr2k( - clblasOrder layout, clblasUplo triangle, clblasTranspose tran_ab, + clblasOrder layout, clblasUplo triangle, clblasTranspose ab_transpose, size_t n, size_t k, float alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, float beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { - return clblasSsyr2k(layout, triangle, tran_ab, + return clblasSsyr2k(layout, triangle, ab_transpose, n, k, alpha, a_mat, a_offset, a_ld, b_mat, b_offset, b_ld, beta, @@ -342,14 +406,14 @@ clblasStatus clblasXsyr2k( num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXsyr2k( - clblasOrder layout, clblasUplo triangle, clblasTranspose tran_ab, + clblasOrder layout, clblasUplo triangle, clblasTranspose ab_transpose, size_t n, size_t k, double alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, double beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { - return clblasDsyr2k(layout, triangle, tran_ab, + return clblasDsyr2k(layout, triangle, ab_transpose, n, k, alpha, a_mat, a_offset, a_ld, b_mat, b_offset, b_ld, beta, @@ -357,16 +421,16 @@ clblasStatus clblasXsyr2k( num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXsyr2k( - clblasOrder layout, clblasUplo triangle, clblasTranspose tran_ab, + clblasOrder layout, clblasUplo triangle, clblasTranspose ab_transpose, size_t n, size_t k, float2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, float2 beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto cl_alpha = cl_float2{{alpha.real(), alpha.imag()}}; auto cl_beta = cl_float2{{beta.real(), beta.imag()}}; - return clblasCsyr2k(layout, triangle, tran_ab, + return clblasCsyr2k(layout, triangle, ab_transpose, n, k, cl_alpha, a_mat, a_offset, a_ld, b_mat, b_offset, b_ld, cl_beta, @@ -374,16 +438,16 @@ clblasStatus clblasXsyr2k( num_queues, queues, num_wait_events, wait_events, events); } clblasStatus clblasXsyr2k( - clblasOrder layout, clblasUplo triangle, clblasTranspose tran_ab, + clblasOrder layout, clblasUplo triangle, clblasTranspose ab_transpose, size_t n, size_t k, double2 alpha, const cl_mem a_mat, size_t a_offset, size_t a_ld, const cl_mem b_mat, size_t b_offset, size_t b_ld, double2 beta, - cl_mem c_mat, size_t c_offset, size_t c_ld, + const cl_mem c_mat, size_t c_offset, size_t c_ld, cl_uint num_queues, cl_command_queue *queues, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto cl_alpha = cl_double2{{alpha.real(), alpha.imag()}}; auto cl_beta = cl_double2{{beta.real(), beta.imag()}}; - return clblasZsyr2k(layout, triangle, tran_ab, + return clblasZsyr2k(layout, triangle, ab_transpose, n, k, cl_alpha, a_mat, a_offset, a_ld, b_mat, b_offset, b_ld, cl_beta, @@ -391,6 +455,160 @@ clblasStatus clblasXsyr2k( num_queues, queues, num_wait_events, wait_events, events); } +// This calls {clblasCher2k, clblasZher2k} with the arguments forwarded. +clblasStatus clblasXher2k( + clblasOrder layout, clblasUplo triangle, clblasTranspose ab_transpose, + size_t n, size_t k, float2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, float beta, + const cl_mem c_mat, size_t c_offset, size_t c_ld, + 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 clblasCher2k(layout, triangle, ab_transpose, + n, k, cl_alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, beta, + c_mat, c_offset, c_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXher2k( + clblasOrder layout, clblasUplo triangle, clblasTranspose ab_transpose, + size_t n, size_t k, double2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, double beta, + const cl_mem c_mat, size_t c_offset, size_t c_ld, + 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 clblasZher2k(layout, triangle, ab_transpose, + n, k, cl_alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, beta, + c_mat, c_offset, c_ld, + num_queues, queues, num_wait_events, wait_events, events); +} + +// This calls {clblasStrmm, clblasDtrmm, clblasCtrmm, clblasZtrmm} with the arguments forwarded. +clblasStatus clblasXtrmm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + clblasTranspose a_transpose, clblasDiag diagonal, + size_t m, size_t n, float alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasStrmm(layout, side, triangle, a_transpose, diagonal, + m, n, alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXtrmm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + clblasTranspose a_transpose, clblasDiag diagonal, + size_t m, size_t n, double alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasDtrmm(layout, side, triangle, a_transpose, diagonal, + m, n, alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXtrmm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + clblasTranspose a_transpose, clblasDiag diagonal, + size_t m, size_t n, float2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, + 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 clblasCtrmm(layout, side, triangle, a_transpose, diagonal, + m, n, cl_alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXtrmm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + clblasTranspose a_transpose, clblasDiag diagonal, + size_t m, size_t n, double2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, + 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 clblasZtrmm(layout, side, triangle, a_transpose, diagonal, + m, n, cl_alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, + num_queues, queues, num_wait_events, wait_events, events); +} + +// This calls {clblasStrsm, clblasDtrsm, clblasCtrsm, clblasZtrsm} with the arguments forwarded. +clblasStatus clblasXtrsm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + clblasTranspose a_transpose, clblasDiag diagonal, + size_t m, size_t n, float alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasStrsm(layout, side, triangle, a_transpose, diagonal, + m, n, alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXtrsm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + clblasTranspose a_transpose, clblasDiag diagonal, + size_t m, size_t n, double alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + return clblasDtrsm(layout, side, triangle, a_transpose, diagonal, + m, n, alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXtrsm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + clblasTranspose a_transpose, clblasDiag diagonal, + size_t m, size_t n, float2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, + 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 clblasCtrsm(layout, side, triangle, a_transpose, diagonal, + m, n, cl_alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, + num_queues, queues, num_wait_events, wait_events, events); +} +clblasStatus clblasXtrsm( + clblasOrder layout, clblasSide side, clblasUplo triangle, + clblasTranspose a_transpose, clblasDiag diagonal, + size_t m, size_t n, double2 alpha, + const cl_mem a_mat, size_t a_offset, size_t a_ld, + const cl_mem b_mat, size_t b_offset, size_t b_ld, + 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 clblasZtrsm(layout, side, triangle, a_transpose, diagonal, + m, n, cl_alpha, + a_mat, a_offset, a_ld, + b_mat, b_offset, b_ld, + num_queues, queues, num_wait_events, wait_events, events); +} + // ================================================================================================= } // namespace clblast |