haraldwolff/CLBlast
1
0
Fork 0
mirror of https://github.com/CNugteren/CLBlast.git synced 2024-07-02 12:26:57 +02:00
Code Issues Packages Projects Releases Wiki Activity

Added the HEMM routine, tester, and client

Browse source
This commit is contained in:
CNugteren 2015-07-12 15:11:50 +02:00
parent 9a929f3fb2
commit b5d39d9d0c
10 changed files with 640 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
CMakeLists.txt
View file

@ -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)
# ==================================================================================================

11
include/clblast.h
View file

@ -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 <typename T>
StatusCode Hemm(const Layout layout, const Side side, const Triangle triangle,
const size_t m, const size_t n,
const T alpha,
const cl_mem a_buffer, const size_t a_offset, const size_t a_ld,
const cl_mem b_buffer, const size_t b_offset, const size_t b_ld,
const T beta,
cl_mem c_buffer, const size_t c_offset, const size_t c_ld,
cl_command_queue* queue, cl_event* event);
// Templated-precision rank-K update of a symmetric matrix: SSYRK/DSYRK/CSYRK/ZSYRK
template <typename T>
StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose a_transpose,

58
include/internal/routines/xhemm.h Normal file
View file

@ -0,0 +1,58 @@
// =================================================================================================
// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This
// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max-
// width of 100 characters per line.
//
// Author(s):
// Cedric Nugteren <www.cedricnugteren.nl>
//
// This file implements the Xhemm routine. It is based on the generalized matrix multiplication
// routine (Xgemm). The implementation is very similar to the Xsymm routine.
//
// =================================================================================================
#ifndef CLBLAST_ROUTINES_XHEMM_H_
#define CLBLAST_ROUTINES_XHEMM_H_
#include "internal/routines/xgemm.h"
namespace clblast {
// =================================================================================================
// See comment at top of file for a description of the class
template <typename T>
class Xhemm: public Xgemm<T> {
public:
// Uses several variables from the Routine class
using Routine::db_;
using Routine::context_;
// Uses several helper functions from the Routine class
using Routine::RunKernel;
using Routine::ErrorIn;
using Routine::TestMatrixA;
using Routine::GetProgramFromCache;
// Uses the regular Xgemm routine
using Xgemm<T>::DoGemm;
// Constructor
Xhemm(CommandQueue &queue, Event &event);
// Templated-precision implementation of the routine
StatusCode DoHemm(const Layout layout, const Side side, const Triangle triangle,
const size_t m, const size_t n,
const T alpha,
const Buffer &a_buffer, const size_t a_offset, const size_t a_ld,
const Buffer &b_buffer, const size_t b_offset, const size_t b_ld,
const T beta,
const Buffer &c_buffer, const size_t c_offset, const size_t c_ld);
};
// =================================================================================================
} // namespace clblast
// CLBLAST_ROUTINES_XHEMM_H_
#endif

49
src/clblast.cc
View file

@ -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<double2>(const Layout, const Side, const Triangle,
// =================================================================================================
// HEMM
template <typename T>
StatusCode Hemm(const Layout layout, const Side side, const Triangle triangle,
const size_t m, const size_t n, const T alpha,
const cl_mem a_buffer, const size_t a_offset, const size_t a_ld,
const cl_mem b_buffer, const size_t b_offset, const size_t b_ld, const T beta,
cl_mem c_buffer, const size_t c_offset, const size_t c_ld,
cl_command_queue* queue, cl_event* event) {
auto queue_cpp = CommandQueue(*queue);
auto event_cpp = Event(*event);
auto routine = Xhemm<T>(queue_cpp, event_cpp);
// Loads the kernel source-code as an include (C++11 raw string literal)
std::string common_source1 =
#include "kernels/copy.opencl"
std::string common_source2 =
#include "kernels/pad.opencl"
std::string common_source3 =
#include "kernels/transpose.opencl"
std::string common_source4 =
#include "kernels/padtranspose.opencl"
std::string kernel_source =
#include "kernels/xgemm.opencl"
auto status = routine.SetUp(common_source1 + common_source2 + common_source3 + common_source4 +
kernel_source);
if (status != StatusCode::kSuccess) { return status; }
// Runs the routine
return routine.DoHemm(layout, side, triangle, m, n, alpha,
Buffer(a_buffer), a_offset, a_ld,
Buffer(b_buffer), b_offset, b_ld, beta,
Buffer(c_buffer), c_offset, c_ld);
}
template StatusCode Hemm<float2>(const Layout, const Side, const Triangle,
const size_t, const size_t, const float2,
const cl_mem, const size_t, const size_t,
const cl_mem, const size_t, const size_t, const float2,
cl_mem, const size_t, const size_t,
cl_command_queue*, cl_event*);
template StatusCode Hemm<double2>(const Layout, const Side, const Triangle,
const size_t, const size_t, const double2,
const cl_mem, const size_t, const size_t,
const cl_mem, const size_t, const size_t, const double2,
cl_mem, const size_t, const size_t,
cl_command_queue*, cl_event*);
// =================================================================================================
// SYRK
template <typename T>
StatusCode Syrk(const Layout layout, const Triangle triangle, const Transpose a_transpose,

83
src/kernels/pad.opencl
View file

@ -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<PAD_WPTX; ++w_one) {
const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0);
#pragma unroll
for (int w_two=0; w_two<PAD_WPTY; ++w_two) {
const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1);
if (id_two < dest_dim && id_one < dest_dim) {
// Loads data from the lower-hermitian matrix
real result;
SetToZero(result);
if (id_two < src_dim && id_one < src_dim) {
if (id_two <= id_one) {
result = src[id_two*src_ld + id_one + src_offset];
if (id_one == id_two) { result.y = ZERO; }
}
else {
result = src[id_one*src_ld + id_two + src_offset];
COMPLEX_CONJUGATE(result);
}
}
// Stores the result in the destination matrix
dest[id_two*dest_ld + id_one + dest_offset] = result;
}
}
}
}
// Same as above, but now the matrix' data is stored in the upper-triangle
__attribute__((reqd_work_group_size(PAD_DIMX, PAD_DIMY, 1)))
__kernel void HermUpperToSquared(const int src_dim,
const int src_ld, const int src_offset,
__global const real* restrict src,
const int dest_dim,
const int dest_ld, const int dest_offset,
__global real* dest) {
// Loops over the work per thread in both dimensions
#pragma unroll
for (int w_one=0; w_one<PAD_WPTX; ++w_one) {
const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0);
#pragma unroll
for (int w_two=0; w_two<PAD_WPTY; ++w_two) {
const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1);
if (id_two < dest_dim && id_one < dest_dim) {
// Loads data from the upper-hermitian matrix
real result;
SetToZero(result);
if (id_two < src_dim && id_one < src_dim) {
if (id_one <= id_two) {
result = src[id_two*src_ld + id_one + src_offset];
if (id_one == id_two) { result.y = ZERO; }
}
else {
result = src[id_one*src_ld + id_two + src_offset];
COMPLEX_CONJUGATE(result);
}
}
// Stores the result in the destination matrix
dest[id_two*dest_ld + id_one + dest_offset] = result;
}
}
}
}
#endif
// =================================================================================================
// Kernel to populate a squared triangular matrix, given that the triangle which holds the data is

130
src/routines/xhemm.cc Normal file
View file

@ -0,0 +1,130 @@
// =================================================================================================
// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This
// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max-
// width of 100 characters per line.
//
// Author(s):
// Cedric Nugteren <www.cedricnugteren.nl>
//
// This file implements the Xhemm class (see the header for information about the class).
//
// =================================================================================================
#include "internal/routines/xhemm.h"
#include <string>
#include <vector>
namespace clblast {
// =================================================================================================
// Constructor: forwards to base class constructor
template <typename T>
Xhemm<T>::Xhemm(CommandQueue &queue, Event &event):
Xgemm<T>(queue, event) {
}
// =================================================================================================
// The main routine
template <typename T>
StatusCode Xhemm<T>::DoHemm(const Layout layout, const Side side, const Triangle triangle,
const size_t m, const size_t n,
const T alpha,
const Buffer &a_buffer, const size_t a_offset, const size_t a_ld,
const Buffer &b_buffer, const size_t b_offset, const size_t b_ld,
const T beta,
const Buffer &c_buffer, const size_t c_offset, const size_t c_ld) {
// Makes sure all dimensions are larger than zero
if ((m == 0) || (n == 0) ) { return StatusCode::kInvalidDimension; }
// Computes the k dimension. This is based on whether or not the hermitian matrix is A (on the
// left) or B (on the right) in the Xgemm routine.
auto k = (side == Side::kLeft) ? m : n;
// Checks for validity of the squared A matrix
auto status = TestMatrixA(k, k, a_buffer, a_offset, a_ld, sizeof(T));
if (ErrorIn(status)) { return status; }
// Determines which kernel to run based on the layout (the Xgemm kernel assumes column-major as
// default) and on whether we are dealing with an upper or lower triangle of the hermitian matrix
bool is_upper = ((triangle == Triangle::kUpper && layout != Layout::kRowMajor) ||
(triangle == Triangle::kLower && layout == Layout::kRowMajor));
auto kernel_name = (is_upper) ? "HermUpperToSquared" : "HermLowerToSquared";
// Temporary buffer for a copy of the hermitian matrix
try {
auto temp_herm = Buffer(context_, CL_MEM_READ_WRITE, k*k*sizeof(T));
// Creates a general matrix from the hermitian matrix to be able to run the regular Xgemm
// routine afterwards
try {
auto& program = GetProgramFromCache();
auto kernel = Kernel(program, kernel_name);
// Sets the arguments for the hermitian-to-squared kernel
kernel.SetArgument(0, static_cast<int>(k));
kernel.SetArgument(1, static_cast<int>(a_ld));
kernel.SetArgument(2, static_cast<int>(a_offset));
kernel.SetArgument(3, a_buffer());
kernel.SetArgument(4, static_cast<int>(k));
kernel.SetArgument(5, static_cast<int>(k));
kernel.SetArgument(6, static_cast<int>(0));
kernel.SetArgument(7, temp_herm());
// Uses the common padding kernel's thread configuration. This is allowed, since the
// hermitian-to-squared kernel uses the same parameters.
auto global = std::vector<size_t>{Ceil(CeilDiv(k, db_["PAD_WPTX"]), db_["PAD_DIMX"]),
Ceil(CeilDiv(k, db_["PAD_WPTY"]), db_["PAD_DIMY"])};
auto local = std::vector<size_t>{db_["PAD_DIMX"], db_["PAD_DIMY"]};
status = RunKernel(kernel, global, local);
if (ErrorIn(status)) { return status; }
// Runs the regular Xgemm code with either "C := AB+C" or ...
if (side == Side::kLeft) {
status = DoGemm(layout, Transpose::kNo, Transpose::kNo,
m, n, k,
alpha,
temp_herm, 0, k,
b_buffer, b_offset, b_ld,
beta,
c_buffer, c_offset, c_ld);
}
// ... with "C := BA+C". Note that A and B are now reversed.
else {
status = DoGemm(layout, Transpose::kNo, Transpose::kNo,
m, n, k,
alpha,
b_buffer, b_offset, b_ld,
temp_herm, 0, k,
beta,
c_buffer, c_offset, c_ld);
// A and B are now reversed, so also reverse the error codes returned from the Xgemm routine
switch(status) {
case StatusCode::kInvalidMatrixA: status = StatusCode::kInvalidMatrixB; break;
case StatusCode::kInvalidMatrixB: status = StatusCode::kInvalidMatrixA; break;
case StatusCode::kInvalidLeadDimA: status = StatusCode::kInvalidLeadDimB; break;
case StatusCode::kInvalidLeadDimB: status = StatusCode::kInvalidLeadDimA; break;
case StatusCode::kInsufficientMemoryA: status = StatusCode::kInsufficientMemoryB; break;
case StatusCode::kInsufficientMemoryB: status = StatusCode::kInsufficientMemoryA; break;
}
}
// Return the status of the Xgemm routine
return status;
} catch (...) { return StatusCode::kInvalidKernel; }
} catch (...) { return StatusCode::kTempBufferAllocFailure; }
}
// =================================================================================================
// Compiles the templated class
template class Xhemm<float2>;
template class Xhemm<double2>;
// =================================================================================================
} // namespace clblast

98
test/correctness/routines/xhemm.cc Normal file
View file

@ -0,0 +1,98 @@
// =================================================================================================
// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This
// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max-
// width of 100 characters per line.
//
// Author(s):
// Cedric Nugteren <www.cedricnugteren.nl>
//
// This file implements the tests for the Xhemm routine.
//
// =================================================================================================
#include "correctness/testblas.h"
#include "routines/xhemm.h"
namespace clblast {
// =================================================================================================
// The correctness tester
template <typename T>
void RunTest(int argc, char *argv[], const bool silent, const std::string &name) {
// Creates a tester
TestBlas<T,T> tester{argc, argv, silent, name, TestXhemm<T>::GetOptions(),
TestXhemm<T>::RunRoutine, TestXhemm<T>::RunReference,
TestXhemm<T>::DownloadResult, TestXhemm<T>::GetResultIndex,
TestXhemm<T>::ResultID1, TestXhemm<T>::ResultID2};
// This variable holds the arguments relevant for this routine
auto args = Arguments<T>{};
// Loops over the test-cases from a data-layout point of view
for (auto &layout: tester.kLayouts) { args.layout = layout;
for (auto &side: tester.kSides) { args.side = side;
for (auto &triangle: tester.kTriangles) { args.triangle = triangle;
// Creates the arguments vector for the regular tests
auto regular_test_vector = std::vector<Arguments<T>>{};
for (auto &m: tester.kMatrixDims) { args.m = m;
for (auto &n: tester.kMatrixDims) { args.n = n;
for (auto &a_ld: tester.kMatrixDims) { args.a_ld = a_ld;
for (auto &a_offset: tester.kOffsets) { args.a_offset = a_offset;
for (auto &b_ld: tester.kMatrixDims) { args.b_ld = b_ld;
for (auto &b_offset: tester.kOffsets) { args.b_offset = b_offset;
for (auto &c_ld: tester.kMatrixDims) { args.c_ld = c_ld;
for (auto &c_offset: tester.kOffsets) { args.c_offset = c_offset;
for (auto &alpha: tester.kAlphaValues) { args.alpha = alpha;
for (auto &beta: tester.kBetaValues) { args.beta = beta;
args.a_size = TestXhemm<T>::GetSizeA(args);
args.b_size = TestXhemm<T>::GetSizeB(args);
args.c_size = TestXhemm<T>::GetSizeC(args);
if (args.a_size<1 || args.b_size<1 || args.c_size<1) { continue; }
regular_test_vector.push_back(args);
}
}
}
}
}
}
}
}
}
}
// Creates the arguments vector for the invalid-buffer tests
auto invalid_test_vector = std::vector<Arguments<T>>{};
args.m = args.n = tester.kBufferSize;
args.a_ld = args.b_ld = args.c_ld = tester.kBufferSize;
args.a_offset = args.b_offset = args.c_offset = 0;
for (auto &a_size: tester.kMatSizes) { args.a_size = a_size;
for (auto &b_size: tester.kMatSizes) { args.b_size = b_size;
for (auto &c_size: tester.kMatSizes) { args.c_size = c_size;
invalid_test_vector.push_back(args);
}
}
}
// Runs the tests
const auto case_name = ToString(layout)+" "+ToString(side)+" "+ToString(triangle);
tester.TestRegular(regular_test_vector, case_name);
tester.TestInvalid(invalid_test_vector, case_name);
}
}
}
}
// =================================================================================================
} // namespace clblast
// Main function (not within the clblast namespace)
int main(int argc, char *argv[]) {
clblast::RunTest<clblast::float2>(argc, argv, true, "CHEMM");
clblast::RunTest<clblast::double2>(argc, argv, true, "ZHEMM");
return 0;
}
// =================================================================================================

40
test/performance/routines/xhemm.cc Normal file
View file

@ -0,0 +1,40 @@
// =================================================================================================
// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This
// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max-
// width of 100 characters per line.
//
// Author(s):
// Cedric Nugteren <www.cedricnugteren.nl>
//
// This file implements the Xhemm command-line interface performance tester.
//
// =================================================================================================
#include "performance/client.h"
#include "routines/xhemm.h"
// =================================================================================================
// Shortcuts to the clblast namespace
using float2 = clblast::float2;
using double2 = clblast::double2;
// Main function (not within the clblast namespace)
int main(int argc, char *argv[]) {
switch(clblast::GetPrecision(argc, argv)) {
case clblast::Precision::kHalf:
throw std::runtime_error("Unsupported precision mode");
case clblast::Precision::kSingle:
throw std::runtime_error("Unsupported precision mode");
case clblast::Precision::kDouble:
throw std::runtime_error("Unsupported precision mode");
case clblast::Precision::kComplexSingle:
clblast::RunClient<clblast::TestXhemm<float2>, float2, float2>(argc, argv); break;
case clblast::Precision::kComplexDouble:
clblast::RunClient<clblast::TestXhemm<double2>, double2, double2>(argc, argv); break;
}
return 0;
}
// =================================================================================================

134
test/routines/xhemm.h Normal file
View file

@ -0,0 +1,134 @@
// =================================================================================================
// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This
// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max-
// width of 100 characters per line.
//
// Author(s):
// Cedric Nugteren <www.cedricnugteren.nl>
//
// This file implements a class with static methods to describe the Xhemm routine. Examples of
// such 'descriptions' are how to calculate the size a of buffer or how to run the routine. These
// static methods are used by the correctness tester and the performance tester.
//
// =================================================================================================
#ifndef CLBLAST_TEST_ROUTINES_XHEMM_H_
#define CLBLAST_TEST_ROUTINES_XHEMM_H_
#include <vector>
#include <string>
#include "wrapper_clblas.h"
namespace clblast {
// =================================================================================================
// See comment at top of file for a description of the class
template <typename T>
class TestXhemm {
public:
// The list of arguments relevant for this routine
static std::vector<std::string> GetOptions() {
return {kArgM, kArgN,
kArgLayout, kArgSide, kArgTriangle,
kArgALeadDim, kArgBLeadDim, kArgCLeadDim,
kArgAOffset, kArgBOffset, kArgCOffset,
kArgAlpha, kArgBeta};
}
// Describes how to obtain the sizes of the buffers
static size_t GetSizeA(const Arguments<T> &args) {
size_t k_value = (args.side == Side::kLeft) ? args.m : args.n;
auto a_rotated = (args.layout == Layout::kRowMajor);
auto a_two = (a_rotated) ? args.m : k_value;
return a_two * args.a_ld + args.a_offset;
}
static size_t GetSizeB(const Arguments<T> &args) {
size_t k_value = (args.side == Side::kLeft) ? args.m : args.n;
auto b_rotated = (args.layout == Layout::kRowMajor);
auto b_two = (b_rotated) ? k_value : args.n;
return b_two * args.b_ld + args.b_offset;
}
static size_t GetSizeC(const Arguments<T> &args) {
auto c_rotated = (args.layout == Layout::kRowMajor);
auto c_two = (c_rotated) ? args.m : args.n;
return c_two * args.c_ld + args.c_offset;
}
// Describes how to set the sizes of all the buffers
static void SetSizes(Arguments<T> &args) {
args.a_size = GetSizeA(args);
args.b_size = GetSizeB(args);
args.c_size = GetSizeC(args);
}
// Describes what the default values of the leading dimensions of the matrices are
static size_t DefaultLDA(const Arguments<T> &args) { return args.m; }
static size_t DefaultLDB(const Arguments<T> &args) { return args.n; }
static size_t DefaultLDC(const Arguments<T> &args) { return args.n; }
// Describes how to run the CLBlast routine
static StatusCode RunRoutine(const Arguments<T> &args, const Buffers &buffers,
CommandQueue &queue) {
auto queue_plain = queue();
auto event = cl_event{};
auto status = Hemm(args.layout, args.side, args.triangle,
args.m, args.n, args.alpha,
buffers.a_mat(), args.a_offset, args.a_ld,
buffers.b_mat(), args.b_offset, args.b_ld, args.beta,
buffers.c_mat(), args.c_offset, args.c_ld,
&queue_plain, &event);
clWaitForEvents(1, &event);
return status;
}
// Describes how to run the clBLAS routine (for correctness/performance comparison)
static StatusCode RunReference(const Arguments<T> &args, const Buffers &buffers,
CommandQueue &queue) {
auto queue_plain = queue();
auto event = cl_event{};
auto status = clblasXhemm(static_cast<clblasOrder>(args.layout),
static_cast<clblasSide>(args.side),
static_cast<clblasUplo>(args.triangle),
args.m, args.n, args.alpha,
buffers.a_mat(), args.a_offset, args.a_ld,
buffers.b_mat(), args.b_offset, args.b_ld, args.beta,
buffers.c_mat(), args.c_offset, args.c_ld,
1, &queue_plain, 0, nullptr, &event);
clWaitForEvents(1, &event);
return static_cast<StatusCode>(status);
}
// Describes how to download the results of the computation (more importantly: which buffer)
static std::vector<T> DownloadResult(const Arguments<T> &args, Buffers &buffers,
CommandQueue &queue) {
std::vector<T> result(args.c_size, static_cast<T>(0));
buffers.c_mat.ReadBuffer(queue, args.c_size*sizeof(T), result);
return result;
}
// Describes how to compute the indices of the result buffer
static size_t ResultID1(const Arguments<T> &args) { return args.m; }
static size_t ResultID2(const Arguments<T> &args) { return args.n; }
static size_t GetResultIndex(const Arguments<T> &args, const size_t id1, const size_t id2) {
return (args.layout == Layout::kRowMajor) ?
id1*args.c_ld + id2 + args.c_offset:
id2*args.c_ld + id1 + args.c_offset;
}
// Describes how to compute performance metrics
static size_t GetFlops(const Arguments<T> &args) {
return 2 * args.m * args.n * args.m;
}
static size_t GetBytes(const Arguments<T> &args) {
return (args.m*args.m + args.m*args.n + 2*args.m*args.n) * sizeof(T);
}
};
// =================================================================================================
} // namespace clblast
// CLBLAST_TEST_ROUTINES_XHEMM_H_
#endif

36
test/wrapper_clblas.h
View file

@ -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,