From 8574f72d46f8f3572e2a5e9f24359d8da18ccf2a Mon Sep 17 00:00:00 2001 From: CNugteren Date: Tue, 30 Jun 2015 07:36:11 +0200 Subject: Added the TRMM and TRSM interface --- include/clblast.h | 29 +++++++++++++++++++++++++---- 1 file changed, 25 insertions(+), 4 deletions(-) (limited to 'include') diff --git a/include/clblast.h b/include/clblast.h index da504a0b..5da10810 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 -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 -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, @@ -131,7 +132,7 @@ StatusCode Symm(const Layout layout, const Side side, const Triangle triangle, // Templated-precision rank-K update of a symmetric matrix: SSYRK/DSYRK/CSYRK/ZSYRK template -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, @@ -141,7 +142,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 -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 +151,26 @@ 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 triangular matrix-matrix multiplication: STRMM/DTRMM/CTRMM/ZTRMM +template +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 +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 -- cgit v1.2.3 From e3dd35f91baf9f0c4cf35d58d3dcbdb4ce8fb3b7 Mon Sep 17 00:00:00 2001 From: CNugteren Date: Wed, 1 Jul 2015 09:39:41 +0200 Subject: Added the unit/non-unit diagonal enum --- include/internal/utilities.h | 2 ++ src/utilities.cc | 8 ++++++++ test/correctness/testblas.h | 1 + test/correctness/tester.cc | 1 + test/performance/client.cc | 2 ++ 5 files changed, 14 insertions(+) (limited to 'include') 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/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(value))+" (unit)"; + case Diagonal::kNonUnit: return ToString(static_cast(value))+" (non-unit)"; + } +} +template <> std::string ToString(Precision value) { switch(value) { case Precision::kHalf: return ToString(static_cast(value))+" (half)"; @@ -143,6 +150,7 @@ template Layout GetArgument(const int, char **, std::string&, const std: template Transpose GetArgument(const int, char **, std::string&, const std::string&, const Transpose); template Side GetArgument(const int, char **, std::string&, const std::string&, const Side); template Triangle GetArgument(const int, char **, std::string&, const std::string&, const Triangle); +template Diagonal GetArgument(const int, char **, std::string&, const std::string&, const Diagonal); template Precision GetArgument(const int, char **, std::string&, const std::string&, const Precision); // ================================================================================================= diff --git a/test/correctness/testblas.h b/test/correctness/testblas.h index 1f92cb30..7469700d 100644 --- a/test/correctness/testblas.h +++ b/test/correctness/testblas.h @@ -61,6 +61,7 @@ class TestBlas: public Tester { const std::vector kLayouts = {Layout::kRowMajor, Layout::kColMajor}; const std::vector kTriangles = {Triangle::kUpper, Triangle::kLower}; const std::vector kSides = {Side::kLeft, Side::kRight}; + const std::vector kDiagonals = {Diagonal::kUnit, Diagonal::kNonUnit}; static const std::vector kTransposes; // Data-type dependent, see .cc-file // Shorthand for the routine-specific functions passed to the tester diff --git a/test/correctness/tester.cc b/test/correctness/tester.cc index 4a179718..db4ee619 100644 --- a/test/correctness/tester.cc +++ b/test/correctness/tester.cc @@ -137,6 +137,7 @@ void Tester::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);} diff --git a/test/performance/client.cc b/test/performance/client.cc index 71471dde..fad0f3a9 100644 --- a/test/performance/client.cc +++ b/test/performance/client.cc @@ -58,6 +58,7 @@ Arguments Client::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}); } @@ -224,6 +225,7 @@ void Client::PrintTableRow(const Arguments& args, const double ms_clblast, else if (o == kArgTriangle) { integers.push_back(static_cast(args.triangle)); } else if (o == kArgATransp) { integers.push_back(static_cast(args.a_transpose)); } else if (o == kArgBTransp) { integers.push_back(static_cast(args.b_transpose)); } + else if (o == kArgDiagonal) { integers.push_back(static_cast(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); } -- cgit v1.2.3 From d9ea0c47c65ff41da2d213cce8b0ef434e817ec2 Mon Sep 17 00:00:00 2001 From: CNugteren Date: Thu, 2 Jul 2015 07:16:04 +0200 Subject: Added the TRMM routine, tester, and client --- CMakeLists.txt | 2 +- include/internal/routines/xtrmm.h | 58 ++++++++++++++++ src/clblast.cc | 4 +- src/kernels/pad.opencl | 97 ++++++++++++++++++++++---- src/routines/xtrmm.cc | 135 +++++++++++++++++++++++++++++++++++++ test/correctness/routines/xtrmm.cc | 96 ++++++++++++++++++++++++++ test/performance/graphs/xtrmm.r | 127 ++++++++++++++++++++++++++++++++++ test/performance/routines/xtrmm.cc | 36 ++++++++++ test/routines/xtrmm.h | 127 ++++++++++++++++++++++++++++++++++ 9 files changed, 666 insertions(+), 16 deletions(-) create mode 100644 include/internal/routines/xtrmm.h create mode 100644 src/routines/xtrmm.cc create mode 100644 test/correctness/routines/xtrmm.cc create mode 100644 test/performance/graphs/xtrmm.r create mode 100644 test/performance/routines/xtrmm.cc create mode 100644 test/routines/xtrmm.h (limited to 'include') diff --git a/CMakeLists.txt b/CMakeLists.txt index 60b1aaed..60060961 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 xsyrk xsyr2k xtrmm) # ================================================================================================== 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 +// +// 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 +class Xtrmm: public Xgemm { + 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::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/src/clblast.cc b/src/clblast.cc index e3ce4d39..299d0a18 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -28,6 +28,7 @@ #include "internal/routines/xsymm.h" #include "internal/routines/xsyrk.h" #include "internal/routines/xsyr2k.h" +#include "internal/routines/xtrmm.h" namespace clblast { // ================================================================================================= @@ -372,7 +373,6 @@ StatusCode Trmm(const Layout layout, const Side side, const Triangle triangle, cl_command_queue* queue, cl_event* event) { auto queue_cpp = CommandQueue(*queue); auto event_cpp = Event(*event); - /* auto routine = Xtrmm(queue_cpp, event_cpp); // Loads the kernel source-code as an include (C++11 raw string literal) @@ -394,8 +394,6 @@ StatusCode Trmm(const Layout layout, const Side side, const Triangle triangle, 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); - */ - return StatusCode::kSuccess; } template StatusCode Trmm(const Layout, const Side, const Triangle, const Transpose, const Diagonal, diff --git a/src/kernels/pad.opencl b/src/kernels/pad.opencl index cce0c746..8294fab7 100644 --- a/src/kernels/pad.opencl +++ b/src/kernels/pad.opencl @@ -135,15 +135,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 +168,88 @@ __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; + } + } + } +} + +// ================================================================================================= + +// 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 +// +// This file implements the Xtrmm class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xtrmm.h" + +#include +#include + +namespace clblast { +// ================================================================================================= + +// Constructor: forwards to base class constructor +template +Xtrmm::Xtrmm(CommandQueue &queue, Event &event): + Xgemm(queue, event) { +} + +// ================================================================================================= + +// The main routine +template +StatusCode Xtrmm::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(k)); + kernel.SetArgument(1, static_cast(a_ld)); + kernel.SetArgument(2, static_cast(a_offset)); + kernel.SetArgument(3, a_buffer()); + kernel.SetArgument(4, static_cast(k)); + kernel.SetArgument(5, static_cast(k)); + kernel.SetArgument(6, static_cast(0)); + kernel.SetArgument(7, temp_triangular()); + kernel.SetArgument(8, static_cast(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{Ceil(CeilDiv(k, db_["PAD_WPTX"]), db_["PAD_DIMX"]), + Ceil(CeilDiv(k, db_["PAD_WPTY"]), db_["PAD_DIMY"])}; + auto local = std::vector{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(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(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; +template class Xtrmm; +template class Xtrmm; +template class Xtrmm; + +// ================================================================================================= +} // namespace clblast diff --git a/test/correctness/routines/xtrmm.cc b/test/correctness/routines/xtrmm.cc new file mode 100644 index 00000000..943fb664 --- /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 +// +// This file implements the tests for the Xtrmm routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/xtrmm.h" + +namespace clblast { +// ================================================================================================= + +// The correctness tester +template +void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { + + // Creates a tester + TestBlas tester{argc, argv, silent, name, TestXtrmm::GetOptions(), + TestXtrmm::RunRoutine, TestXtrmm::RunReference, + TestXtrmm::DownloadResult, TestXtrmm::GetResultIndex, + TestXtrmm::ResultID1, TestXtrmm::ResultID2}; + + // This variable holds the arguments relevant for this routine + auto args = Arguments{}; + + // 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>{}; + 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::GetSizeA(args); + args.b_size = TestXtrmm::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>{}; + 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(argc, argv, false, "STRMM"); + clblast::RunTest(argc, argv, true, "DTRMM"); + clblast::RunTest(argc, argv, true, "CTRMM"); + clblast::RunTest(argc, argv, true, "ZTRMM"); + return 0; +} + +// ================================================================================================= 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 +# +# 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/xtrmm.cc b/test/performance/routines/xtrmm.cc new file mode 100644 index 00000000..91dcbd07 --- /dev/null +++ b/test/performance/routines/xtrmm.cc @@ -0,0 +1,36 @@ + +// ================================================================================================= +// 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 +// +// This file implements the Xtrmm command-line interface performance tester. +// +// ================================================================================================= + +#include "performance/client.h" +#include "routines/xtrmm.h" + +// ================================================================================================= + +// 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, float>(argc, argv); break; + case clblast::Precision::kDouble: + clblast::RunClient, double>(argc, argv); break; + case clblast::Precision::kComplexSingle: + clblast::RunClient, clblast::float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient, clblast::double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= 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 +// +// 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 +#include + +#include "wrapper_clblas.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template +class TestXtrmm { + public: + + // The list of arguments relevant for this routine + static std::vector 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 &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 &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 &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 &args) { return args.m; } + static size_t DefaultLDB(const Arguments &args) { return args.n; } + static size_t DefaultLDC(const Arguments &) { return 1; } // N/A for this routine + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments &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 &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXtrmm(static_cast(args.layout), + static_cast(args.side), + static_cast(args.triangle), + static_cast(args.a_transpose), + static_cast(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(status); + } + + // Describes how to download the results of the computation (more importantly: which buffer) + static std::vector DownloadResult(const Arguments &args, Buffers &buffers, + CommandQueue &queue) { + std::vector result(args.b_size, static_cast(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 &args) { return args.m; } + static size_t ResultID2(const Arguments &args) { return args.n; } + static size_t GetResultIndex(const Arguments &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 &args) { + auto k = (args.side == Side::kLeft) ? args.m : args.n; + return args.m * args.n * k; + } + static size_t GetBytes(const Arguments &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 -- cgit v1.2.3 From 5578d5ab282d63ad47a767dcbebb94b83195230d Mon Sep 17 00:00:00 2001 From: CNugteren Date: Wed, 8 Jul 2015 07:25:18 +0200 Subject: Added option to set the imaginary part of the diagonal to zero --- include/internal/routine.h | 1 + src/routine.cc | 4 +++- src/routines/xgemm.cc | 8 ++++---- src/routines/xsyr2k.cc | 11 ++++------- src/routines/xsyrk.cc | 11 ++++------- 5 files changed, 16 insertions(+), 19 deletions(-) (limited to 'include') 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/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(upper)); kernel.SetArgument(11, static_cast(lower)); + kernel.SetArgument(12, static_cast(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::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(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::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/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::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::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::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::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::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::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 -- cgit v1.2.3 From 919bba3eaf0feaa83e787aa500d6f0d5169b02b5 Mon Sep 17 00:00:00 2001 From: CNugteren Date: Fri, 10 Jul 2015 07:19:59 +0200 Subject: Added the HERK routine, tester, and client --- CMakeLists.txt | 2 +- include/clblast.h | 10 +++ include/internal/routines/xherk.h | 47 +++++++++++ src/clblast.cc | 47 ++++++++++- src/routines/xherk.cc | 156 +++++++++++++++++++++++++++++++++++++ test/performance/routines/xherk.cc | 40 ++++++++++ test/routines/xherk.h | 121 ++++++++++++++++++++++++++++ test/wrapper_clblas.h | 28 +++++++ 8 files changed, 449 insertions(+), 2 deletions(-) create mode 100644 include/internal/routines/xherk.h create mode 100644 src/routines/xherk.cc create mode 100644 test/performance/routines/xherk.cc create mode 100644 test/routines/xherk.h (limited to 'include') diff --git a/CMakeLists.txt b/CMakeLists.txt index 60060961..61af6b82 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -98,7 +98,7 @@ set(SAMPLE_PROGRAMS sgemm) set(ROUTINES xaxpy xgemv - xgemm xsymm xsyrk xsyr2k xtrmm) + xgemm xsymm xsyrk xherk xsyr2k xtrmm) # ================================================================================================== diff --git a/include/clblast.h b/include/clblast.h index 5da10810..d7c902d9 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -140,6 +140,16 @@ StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose a_ 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 +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 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-2K update of a symmetric matrix: SSYR2K/DSYR2K/CSYR2K/ZSYR2K template StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose ab_transpose, 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 +// +// 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 +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/src/clblast.cc b/src/clblast.cc index 299d0a18..638bc944 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -27,6 +27,7 @@ #include "internal/routines/xgemm.h" #include "internal/routines/xsymm.h" #include "internal/routines/xsyrk.h" +#include "internal/routines/xherk.h" #include "internal/routines/xsyr2k.h" #include "internal/routines/xtrmm.h" @@ -302,6 +303,50 @@ template StatusCode Syrk(const Layout, const Triangle, const Transpose, // ================================================================================================= +// HERK +template +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 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,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(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(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 StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose ab_transpose, @@ -456,7 +501,7 @@ StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, Buffer(a_buffer), a_offset, a_ld, Buffer(b_buffer), b_offset, b_ld); */ - return StatusCode::kSuccess; + return StatusCode::kNotImplemented; } template StatusCode Trsm(const Layout, const Side, const Triangle, const Transpose, const Diagonal, 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 +// +// This file implements the Xherk class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xherk.h" + +#include +#include + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xherk::precision_ = Precision::kComplexSingle; +template <> const Precision Xherk::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template +Xherk::Xherk(CommandQueue &queue, Event &event): + Routine(queue, event, {"Copy", "Pad", "Transpose", "PadTranspose", "Xgemm"}, precision_) { +} + +// ================================================================================================= + +// The main routine +template +StatusCode Xherk::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(0.0)}; + auto complex_beta = T{beta, static_cast(0.0)}; + kernel.SetArgument(0, static_cast(n_ceiled)); + kernel.SetArgument(1, static_cast(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{ + (n_ceiled * db_["MDIMC"]) / db_["MWG"], + (n_ceiled * db_["NDIMC"]) / db_["NWG"] + }; + auto local = std::vector{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; +template class Xherk; + +// ================================================================================================= +} // namespace clblast 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 +// +// 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, float2, float>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient, double2, double>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= 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 +// +// 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 +#include + +#include "wrapper_clblas.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template +class TestXherk { + public: + + // The list of arguments relevant for this routine + static std::vector 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 &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 &args) { + return args.n * args.c_ld + args.c_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments &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 &args) { return args.k; } + static size_t DefaultLDB(const Arguments &) { return 1; } // N/A for this routine + static size_t DefaultLDC(const Arguments &args) { return args.n; } + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments &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 &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXherk(static_cast(args.layout), + static_cast(args.triangle), + static_cast(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(status); + } + + // Describes how to download the results of the computation (more importantly: which buffer) + static std::vector DownloadResult(const Arguments &args, Buffers &buffers, + CommandQueue &queue) { + std::vector result(args.c_size, static_cast(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 &args) { return args.n; } + static size_t ResultID2(const Arguments &args) { return args.n; } + static size_t GetResultIndex(const Arguments &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 &args) { + return args.n * args.n * args.k; + } + static size_t GetBytes(const Arguments &args) { + return (args.n*args.k + args.n*args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XHERK_H_ +#endif diff --git a/test/wrapper_clblas.h b/test/wrapper_clblas.h index a03ade5e..360adfc8 100644 --- a/test/wrapper_clblas.h +++ b/test/wrapper_clblas.h @@ -325,6 +325,34 @@ clblasStatus clblasXsyrk( 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 ab_transpose, -- cgit v1.2.3 From b02876d6e9f711369474219576e7bcbebdb10e1c Mon Sep 17 00:00:00 2001 From: CNugteren Date: Fri, 10 Jul 2015 20:59:20 +0200 Subject: Added the HER2K routine, tester, and client --- CMakeLists.txt | 2 +- include/clblast.h | 11 +++ include/internal/routines/xher2k.h | 48 ++++++++++ src/clblast.cc | 59 +++++++++++- src/routines/xher2k.cc | 178 ++++++++++++++++++++++++++++++++++++ test/correctness/routines/xher2k.cc | 100 ++++++++++++++++++++ test/performance/routines/xher2k.cc | 40 ++++++++ test/routines/xher2k.h | 132 ++++++++++++++++++++++++++ test/wrapper_clblas.h | 34 +++++++ 9 files changed, 598 insertions(+), 6 deletions(-) create mode 100644 include/internal/routines/xher2k.h create mode 100644 src/routines/xher2k.cc create mode 100644 test/correctness/routines/xher2k.cc create mode 100644 test/performance/routines/xher2k.cc create mode 100644 test/routines/xher2k.h (limited to 'include') diff --git a/CMakeLists.txt b/CMakeLists.txt index 61af6b82..96e6573e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -98,7 +98,7 @@ set(SAMPLE_PROGRAMS sgemm) set(ROUTINES xaxpy xgemv - xgemm xsymm xsyrk xherk xsyr2k xtrmm) + xgemm xsymm xsyrk xherk xsyr2k xher2k xtrmm) # ================================================================================================== diff --git a/include/clblast.h b/include/clblast.h index d7c902d9..ef279fe5 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -161,6 +161,17 @@ StatusCode Syr2k(const Layout layout, const Triangle triangle, const Transpose a 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 +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 StatusCode Trmm(const Layout layout, const Side side, const Triangle triangle, 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 +// +// 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 +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/src/clblast.cc b/src/clblast.cc index 638bc944..00a90707 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -29,6 +29,7 @@ #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 { @@ -350,11 +351,11 @@ template StatusCode Herk(const Layout, const Triangle, const Transpose, // SYR2K template 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) { + 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(queue_cpp, event_cpp); @@ -407,6 +408,54 @@ template StatusCode Syr2k(const Layout, const Triangle, const Transpose // ================================================================================================= +// SYR2K +template +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(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(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(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 StatusCode Trmm(const Layout layout, const Side side, const Triangle triangle, 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 +// +// This file implements the Xher2k class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xher2k.h" + +#include +#include + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xher2k::precision_ = Precision::kComplexSingle; +template <> const Precision Xher2k::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template +Xher2k::Xher2k(CommandQueue &queue, Event &event): + Routine(queue, event, {"Copy", "Pad", "Transpose", "PadTranspose", "Xgemm"}, precision_) { +} + +// ================================================================================================= + +// The main routine +template +StatusCode Xher2k::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(0.0)}; + kernel.SetArgument(0, static_cast(n_ceiled)); + kernel.SetArgument(1, static_cast(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{ + (n_ceiled * db_["MDIMC"]) / db_["MWG"], + (n_ceiled * db_["NDIMC"]) / db_["NWG"] + }; + auto local = std::vector{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(1.0), static_cast(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; +template class Xher2k; + +// ================================================================================================= +} // namespace clblast 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 +// +// This file implements the tests for the Xher2k routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/xher2k.h" + +namespace clblast { +// ================================================================================================= + +// The correctness tester +template +void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { + + // Creates a tester + TestBlas tester{argc, argv, silent, name, TestXher2k::GetOptions(), + TestXher2k::RunRoutine, TestXher2k::RunReference, + TestXher2k::DownloadResult, TestXher2k::GetResultIndex, + TestXher2k::ResultID1, TestXher2k::ResultID2}; + + // This variable holds the arguments relevant for this routine + auto args = Arguments{}; + + // 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>{}; + 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::GetSizeA(args); + args.b_size = TestXher2k::GetSizeB(args); + args.c_size = TestXher2k::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>{}; + 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(argc, argv, false, "CHER2K"); + clblast::RunTest(argc, argv, true, "ZHER2K"); + 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 +// +// 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, float2, float>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient, double2, double>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= 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 +// +// 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 +#include + +#include "wrapper_clblas.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template +class TestXher2k { + public: + + // The list of arguments relevant for this routine + static std::vector 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 &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 &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 &args) { + return args.n * args.c_ld + args.c_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments &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 &args) { return args.k; } + static size_t DefaultLDB(const Arguments &args) { return args.k; } + static size_t DefaultLDC(const Arguments &args) { return args.n; } + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments &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 &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(args.layout), + static_cast(args.triangle), + static_cast(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(status); + } + + // Describes how to download the results of the computation (more importantly: which buffer) + static std::vector DownloadResult(const Arguments &args, Buffers &buffers, + CommandQueue &queue) { + std::vector result(args.c_size, static_cast(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 &args) { return args.n; } + static size_t ResultID2(const Arguments &args) { return args.n; } + static size_t GetResultIndex(const Arguments &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 &args) { + return 2 * args.n * args.n * args.k; + } + static size_t GetBytes(const Arguments &args) { + return (args.n*args.k + args.n*args.n) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XHER2K_H_ +#endif diff --git a/test/wrapper_clblas.h b/test/wrapper_clblas.h index 360adfc8..03f31a58 100644 --- a/test/wrapper_clblas.h +++ b/test/wrapper_clblas.h @@ -419,6 +419,40 @@ 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, -- cgit v1.2.3 From 9a929f3fb2081bd2fd8f68efce3b9d93e86bf611 Mon Sep 17 00:00:00 2001 From: CNugteren Date: Fri, 10 Jul 2015 21:08:18 +0200 Subject: Disabled prototype of TRSM --- include/clblast.h | 2 ++ src/clblast.cc | 6 ++---- 2 files changed, 4 insertions(+), 4 deletions(-) (limited to 'include') diff --git a/include/clblast.h b/include/clblast.h index ef279fe5..e6c49dbf 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -183,6 +183,7 @@ StatusCode Trmm(const Layout layout, const Side side, const Triangle triangle, cl_command_queue* queue, cl_event* event); // Templated-precision matrix equation solver: STRSM/DTRSM/CTRSM/ZTRSM +/* template StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, const Transpose a_transpose, const Diagonal diagonal, @@ -191,6 +192,7 @@ StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, 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/src/clblast.cc b/src/clblast.cc index 00a90707..66202adb 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -517,6 +517,7 @@ template StatusCode Trmm(const Layout, const Side, const Triangle, // ================================================================================================= // TRSM +/* template StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, const Transpose a_transpose, const Diagonal diagonal, @@ -527,7 +528,6 @@ StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, cl_command_queue* queue, cl_event* event) { auto queue_cpp = CommandQueue(*queue); auto event_cpp = Event(*event); - /* auto routine = Xtrsm(queue_cpp, event_cpp); // Loads the kernel source-code as an include (C++11 raw string literal) @@ -549,8 +549,6 @@ StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, 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); - */ - return StatusCode::kNotImplemented; } template StatusCode Trsm(const Layout, const Side, const Triangle, const Transpose, const Diagonal, @@ -576,6 +574,6 @@ template StatusCode Trsm(const Layout, const Side, const Triangle, const cl_mem, const size_t, const size_t, cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); - +*/ // ================================================================================================= } // namespace clblast -- cgit v1.2.3 From b5d39d9d0c3e1084cb5131e2822d4fb754b0b412 Mon Sep 17 00:00:00 2001 From: CNugteren Date: Sun, 12 Jul 2015 15:11:50 +0200 Subject: Added the HEMM routine, tester, and client --- CMakeLists.txt | 2 +- include/clblast.h | 11 +++ include/internal/routines/xhemm.h | 58 ++++++++++++++++ src/clblast.cc | 49 ++++++++++++++ src/kernels/pad.opencl | 83 +++++++++++++++++++++++ src/routines/xhemm.cc | 130 +++++++++++++++++++++++++++++++++++ test/correctness/routines/xhemm.cc | 98 +++++++++++++++++++++++++++ test/performance/routines/xhemm.cc | 40 +++++++++++ test/routines/xhemm.h | 134 +++++++++++++++++++++++++++++++++++++ test/wrapper_clblas.h | 36 ++++++++++ 10 files changed, 640 insertions(+), 1 deletion(-) create mode 100644 include/internal/routines/xhemm.h create mode 100644 src/routines/xhemm.cc create mode 100644 test/correctness/routines/xhemm.cc create mode 100644 test/performance/routines/xhemm.cc create mode 100644 test/routines/xhemm.h (limited to 'include') diff --git a/CMakeLists.txt b/CMakeLists.txt index 96e6573e..1b2c5657 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -98,7 +98,7 @@ set(SAMPLE_PROGRAMS sgemm) set(ROUTINES xaxpy xgemv - xgemm xsymm xsyrk xherk xsyr2k xher2k xtrmm) + xgemm xsymm xhemm xsyrk xherk xsyr2k xher2k xtrmm) # ================================================================================================== diff --git a/include/clblast.h b/include/clblast.h index e6c49dbf..80ea1707 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -130,6 +130,17 @@ 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 +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 StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose a_transpose, 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 +// +// 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 +class Xhemm: public Xgemm { + 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::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/src/clblast.cc b/src/clblast.cc index 66202adb..23046b01 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -26,6 +26,7 @@ // 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" @@ -250,6 +251,54 @@ template StatusCode Symm(const Layout, const Side, const Triangle, // ================================================================================================= +// HEMM +template +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(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(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(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 StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose a_transpose, diff --git a/src/kernels/pad.opencl b/src/kernels/pad.opencl index f8a89d24..2791db30 100644 --- a/src/kernels/pad.opencl +++ b/src/kernels/pad.opencl @@ -185,6 +185,89 @@ __kernel void SymmUpperToSquared(const int src_dim, } } +// ================================================================================================= +#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 +// +// This file implements the Xhemm class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/xhemm.h" + +#include +#include + +namespace clblast { +// ================================================================================================= + +// Constructor: forwards to base class constructor +template +Xhemm::Xhemm(CommandQueue &queue, Event &event): + Xgemm(queue, event) { +} + +// ================================================================================================= + +// The main routine +template +StatusCode Xhemm::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(k)); + kernel.SetArgument(1, static_cast(a_ld)); + kernel.SetArgument(2, static_cast(a_offset)); + kernel.SetArgument(3, a_buffer()); + kernel.SetArgument(4, static_cast(k)); + kernel.SetArgument(5, static_cast(k)); + kernel.SetArgument(6, static_cast(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{Ceil(CeilDiv(k, db_["PAD_WPTX"]), db_["PAD_DIMX"]), + Ceil(CeilDiv(k, db_["PAD_WPTY"]), db_["PAD_DIMY"])}; + auto local = std::vector{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; +template class Xhemm; + +// ================================================================================================= +} // namespace clblast 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 +// +// This file implements the tests for the Xhemm routine. +// +// ================================================================================================= + +#include "correctness/testblas.h" +#include "routines/xhemm.h" + +namespace clblast { +// ================================================================================================= + +// The correctness tester +template +void RunTest(int argc, char *argv[], const bool silent, const std::string &name) { + + // Creates a tester + TestBlas tester{argc, argv, silent, name, TestXhemm::GetOptions(), + TestXhemm::RunRoutine, TestXhemm::RunReference, + TestXhemm::DownloadResult, TestXhemm::GetResultIndex, + TestXhemm::ResultID1, TestXhemm::ResultID2}; + + // This variable holds the arguments relevant for this routine + auto args = Arguments{}; + + // 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>{}; + 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::GetSizeA(args); + args.b_size = TestXhemm::GetSizeB(args); + args.c_size = TestXhemm::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>{}; + 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(argc, argv, true, "CHEMM"); + clblast::RunTest(argc, argv, true, "ZHEMM"); + 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 +// +// 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, float2, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: + clblast::RunClient, 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 +// +// 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 +#include + +#include "wrapper_clblas.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template +class TestXhemm { + public: + + // The list of arguments relevant for this routine + static std::vector 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 &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 &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 &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 &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 &args) { return args.m; } + static size_t DefaultLDB(const Arguments &args) { return args.n; } + static size_t DefaultLDC(const Arguments &args) { return args.n; } + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments &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 &args, const Buffers &buffers, + CommandQueue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXhemm(static_cast(args.layout), + static_cast(args.side), + static_cast(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(status); + } + + // Describes how to download the results of the computation (more importantly: which buffer) + static std::vector DownloadResult(const Arguments &args, Buffers &buffers, + CommandQueue &queue) { + std::vector result(args.c_size, static_cast(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 &args) { return args.m; } + static size_t ResultID2(const Arguments &args) { return args.n; } + static size_t GetResultIndex(const Arguments &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 &args) { + return 2 * args.m * args.n * args.m; + } + static size_t GetBytes(const Arguments &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/wrapper_clblas.h b/test/wrapper_clblas.h index 03f31a58..4aaf3705 100644 --- a/test/wrapper_clblas.h +++ b/test/wrapper_clblas.h @@ -267,6 +267,42 @@ 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 a_transpose, -- cgit v1.2.3