From a2e726d3bd4294f1eae1735f6ba23105dccc6b10 Mon Sep 17 00:00:00 2001 From: CNugteren Date: Mon, 14 Sep 2015 16:57:00 +0200 Subject: Added xDOT/xDOTU/xDOTC dot-product routines --- src/clblast.cc | 87 ++++++++++++++++++++++++++++ src/clblast_c.cc | 78 +++++++++++++++++++++++++ src/database.cc | 2 + src/kernels/common.opencl | 7 +++ src/kernels/level1/xdot.opencl | 106 ++++++++++++++++++++++++++++++++++ src/routines/level1/xdot.cc | 115 +++++++++++++++++++++++++++++++++++++ src/routines/level1/xdotc.cc | 49 ++++++++++++++++ src/routines/level1/xdotu.cc | 49 ++++++++++++++++ src/tuning/xdot.cc | 125 +++++++++++++++++++++++++++++++++++++++++ 9 files changed, 618 insertions(+) create mode 100644 src/kernels/level1/xdot.opencl create mode 100644 src/routines/level1/xdot.cc create mode 100644 src/routines/level1/xdotc.cc create mode 100644 src/routines/level1/xdotu.cc create mode 100644 src/tuning/xdot.cc (limited to 'src') diff --git a/src/clblast.cc b/src/clblast.cc index c99ad7b1..3303085e 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -22,6 +22,9 @@ #include "internal/routines/level1/xscal.h" #include "internal/routines/level1/xcopy.h" #include "internal/routines/level1/xaxpy.h" +#include "internal/routines/level1/xdot.h" +#include "internal/routines/level1/xdotu.h" +#include "internal/routines/level1/xdotc.h" // BLAS level-2 includes #include "internal/routines/level2/xgemv.h" @@ -177,6 +180,90 @@ template StatusCode Axpy(const size_t, cl_mem, const size_t, const size_t, cl_command_queue* queue, cl_event* event); +// DOT +template +StatusCode Dot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xdot(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoDot(n, + Buffer(dot_buffer), dot_offset, + Buffer(x_buffer), x_offset, x_inc, + Buffer(y_buffer), y_offset, y_inc); +} +template StatusCode Dot(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Dot(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); + +// DOTU +template +StatusCode Dotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xdotu(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoDotu(n, + Buffer(dot_buffer), dot_offset, + Buffer(x_buffer), x_offset, x_inc, + Buffer(y_buffer), y_offset, y_inc); +} +template StatusCode Dotu(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Dotu(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); + +// DOTC +template +StatusCode Dotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xdotc(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoDotc(n, + Buffer(dot_buffer), dot_offset, + Buffer(x_buffer), x_offset, x_inc, + Buffer(y_buffer), y_offset, y_inc); +} +template StatusCode Dotc(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); +template StatusCode Dotc(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue* queue, cl_event* event); + // ================================================================================================= // BLAS level-2 (matrix-vector) routines // ================================================================================================= diff --git a/src/clblast_c.cc b/src/clblast_c.cc index ab3994fb..eccf517f 100644 --- a/src/clblast_c.cc +++ b/src/clblast_c.cc @@ -203,6 +203,84 @@ StatusCode CLBlastZaxpy(const size_t n, return static_cast(status); } +// DOT +StatusCode CLBlastSdot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dot(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast(status); +} +StatusCode CLBlastDdot(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dot(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast(status); +} + +// DOTU +StatusCode CLBlastCdotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dotu(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast(status); +} +StatusCode CLBlastZdotu(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dotu(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast(status); +} + +// DOTC +StatusCode CLBlastCdotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dotc(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast(status); +} +StatusCode CLBlastZdotc(const size_t n, + cl_mem dot_buffer, const size_t dot_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Dotc(n, + dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + queue, event); + return static_cast(status); +} + // ================================================================================================= // BLAS level-2 (matrix-vector) routines // ================================================================================================= diff --git a/src/database.cc b/src/database.cc index 258d861e..b7275dad 100644 --- a/src/database.cc +++ b/src/database.cc @@ -13,6 +13,7 @@ #include "internal/database.h" #include "internal/database/xaxpy.h" +#include "internal/database/xdot.h" #include "internal/database/xgemv.h" #include "internal/database/xgemm.h" #include "internal/database/copy.h" @@ -28,6 +29,7 @@ namespace clblast { // Initializes the database const std::vector Database::database = { XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble, + XdotSingle, XdotDouble, XdotComplexSingle, XdotComplexDouble, XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble, XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble, CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble, diff --git a/src/kernels/common.opencl b/src/kernels/common.opencl index 8e71429e..f2a2e7a7 100644 --- a/src/kernels/common.opencl +++ b/src/kernels/common.opencl @@ -109,6 +109,13 @@ R"( #define SetToOne(a) a = ONE #endif +// Adds two complex variables +#if PRECISION == 3232 || PRECISION == 6464 + #define Add(c, a, b) c.x = a.x + b.x; c.y = a.y + b.y +#else + #define Add(c, a, b) c = a + b +#endif + // Multiply two complex variables (used in the defines below) #if PRECISION == 3232 || PRECISION == 6464 #define MulReal(a, b) a.x*b.x - a.y*b.y diff --git a/src/kernels/level1/xdot.opencl b/src/kernels/level1/xdot.opencl new file mode 100644 index 00000000..e13eb3c1 --- /dev/null +++ b/src/kernels/level1/xdot.opencl @@ -0,0 +1,106 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren +// +// This file contains the Xdot kernel. It implements a dot-product computation using reduction +// kernels. Reduction is split in two parts. In the first (main) kernel the X and Y vectors are +// multiplied, followed by a per-thread and a per-workgroup reduction. The second (epilogue) kernel +// is executed with a single workgroup only, computing the final result. +// +// ================================================================================================= + +// Enables loading of this file using the C++ pre-processor's #include (C++11 standard raw string +// literal). Comment-out this line for syntax-highlighting when developing. +R"( + +// Parameters set by the tuner or by the database. Here they are given a basic default value in case +// this kernel file is used outside of the CLBlast library. +#ifndef WGS1 + #define WGS1 64 // The local work-group size of the main kernel +#endif +#ifndef WGS2 + #define WGS2 64 // The local work-group size of the epilogue kernel +#endif + +// ================================================================================================= + +// The main reduction kernel, performing the multiplication and the majority of the sum operation +__attribute__((reqd_work_group_size(WGS1, 1, 1))) +__kernel void Xdot(const int n, + const __global real* restrict xgm, const int x_offset, const int x_inc, + const __global real* restrict ygm, const int y_offset, const int y_inc, + __global real* output, const int do_conjugate) { + __local real lm[WGS1]; + const int lid = get_local_id(0); + const int wgid = get_group_id(0); + const int num_groups = get_num_groups(0); + + // Performs multiplication and the first steps of the reduction + real acc; + SetToZero(acc); + int id = wgid*WGS1 + lid; + while (id < n) { + real x = xgm[id*x_inc + x_offset]; + real y = ygm[id*y_inc + y_offset]; + if (do_conjugate) { COMPLEX_CONJUGATE(x); } + MultiplyAdd(acc, x, y); + id += WGS1*num_groups; + } + lm[lid] = acc; + barrier(CLK_LOCAL_MEM_FENCE); + + // Performs reduction in local memory + #pragma unroll + for (int s=WGS1/2; s>0; s=s>>1) { + if (lid < s) { + Add(lm[lid], lm[lid], lm[lid + s]); + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + // Stores the per-workgroup result + if (lid == 0) { + output[wgid] = lm[0]; + } +} + +// ================================================================================================= + +// The epilogue reduction kernel, performing the final bit of the sum operation. This kernel has to +// be launched with a single workgroup only. +__attribute__((reqd_work_group_size(WGS2, 1, 1))) +__kernel void XdotEpilogue(const __global real* restrict input, + __global real* dot, const int dot_offset) { + __local real lm[WGS2]; + const int lid = get_local_id(0); + + // Performs the first step of the reduction while loading the data + Add(lm[lid], input[lid], input[lid + WGS2]); + barrier(CLK_LOCAL_MEM_FENCE); + + // Performs reduction in local memory + #pragma unroll + for (int s=WGS2/2; s>0; s=s>>1) { + if (lid < s) { + Add(lm[lid], lm[lid], lm[lid + s]); + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + // Stores the final result + if (lid == 0) { + dot[dot_offset] = lm[0]; + } +} + +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= diff --git a/src/routines/level1/xdot.cc b/src/routines/level1/xdot.cc new file mode 100644 index 00000000..a0c1e756 --- /dev/null +++ b/src/routines/level1/xdot.cc @@ -0,0 +1,115 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren +// +// This file implements the Xdot class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xdot.h" + +#include +#include + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xdot::precision_ = Precision::kSingle; +template <> const Precision Xdot::precision_ = Precision::kDouble; +template <> const Precision Xdot::precision_ = Precision::kComplexSingle; +template <> const Precision Xdot::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template +Xdot::Xdot(Queue &queue, Event &event, const std::string &name): + Routine(queue, event, name, {"Xdot"}, precision_) { + source_string_ = + #include "../../kernels/level1/xdot.opencl" + ; +} + +// ================================================================================================= + +// The main routine +template +StatusCode Xdot::DoDot(const size_t n, + const Buffer &dot_buffer, const size_t dot_offset, + const Buffer &x_buffer, const size_t x_offset, const size_t x_inc, + const Buffer &y_buffer, const size_t y_offset, const size_t y_inc, + const bool do_conjugate) { + + // Makes sure all dimensions are larger than zero + if (n == 0) { return StatusCode::kInvalidDimension; } + + // Tests the vectors for validity + auto status = TestVectorX(n, x_buffer, x_offset, x_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestVectorY(n, y_buffer, y_offset, y_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestVectorDot(1, dot_buffer, dot_offset, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Retrieves the Xdot kernels from the compiled binary + try { + auto& program = GetProgramFromCache(); + auto kernel1 = Kernel(program, "Xdot"); + auto kernel2 = Kernel(program, "XdotEpilogue"); + + // Creates the buffer for intermediate values + auto temp_size = 2*db_["WGS2"]; + auto temp_buffer = Buffer(context_, temp_size); + + // Sets the kernel arguments + kernel1.SetArgument(0, static_cast(n)); + kernel1.SetArgument(1, x_buffer()); + kernel1.SetArgument(2, static_cast(x_offset)); + kernel1.SetArgument(3, static_cast(x_inc)); + kernel1.SetArgument(4, y_buffer()); + kernel1.SetArgument(5, static_cast(y_offset)); + kernel1.SetArgument(6, static_cast(y_inc)); + kernel1.SetArgument(7, temp_buffer()); + kernel1.SetArgument(8, static_cast(do_conjugate)); + + // Launches the main kernel + auto global1 = std::vector{db_["WGS1"]*temp_size}; + auto local1 = std::vector{db_["WGS1"]}; + status = RunKernel(kernel1, global1, local1); + if (ErrorIn(status)) { return status; } + + // Sets the arguments for the epilogue kernel + kernel2.SetArgument(0, temp_buffer()); + kernel2.SetArgument(1, dot_buffer()); + kernel2.SetArgument(2, static_cast(dot_offset)); + + // Launches the epilogue kernel + auto global2 = std::vector{db_["WGS2"]}; + auto local2 = std::vector{db_["WGS2"]}; + status = RunKernel(kernel2, global2, local2); + if (ErrorIn(status)) { return status; } + + // Waits for all kernels to finish + queue_.Finish(); + + // Succesfully finished the computation + return StatusCode::kSuccess; + } catch (...) { return StatusCode::kInvalidKernel; } +} + +// ================================================================================================= + +// Compiles the templated class +template class Xdot; +template class Xdot; +template class Xdot; +template class Xdot; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/level1/xdotc.cc b/src/routines/level1/xdotc.cc new file mode 100644 index 00000000..f414f556 --- /dev/null +++ b/src/routines/level1/xdotc.cc @@ -0,0 +1,49 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren +// +// This file implements the Xdotc class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xdotc.h" + +#include +#include + +namespace clblast { +// ================================================================================================= + +// Constructor: forwards to base class constructor +template +Xdotc::Xdotc(Queue &queue, Event &event, const std::string &name): + Xdot(queue, event, name) { +} + +// ================================================================================================= + +// The main routine +template +StatusCode Xdotc::DoDotc(const size_t n, + const Buffer &dot_buffer, const size_t dot_offset, + const Buffer &x_buffer, const size_t x_offset, const size_t x_inc, + const Buffer &y_buffer, const size_t y_offset, const size_t y_inc) { + return DoDot(n, dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + true); +} + +// ================================================================================================= + +// Compiles the templated class +template class Xdotc; +template class Xdotc; + +// ================================================================================================= +} // namespace clblast diff --git a/src/routines/level1/xdotu.cc b/src/routines/level1/xdotu.cc new file mode 100644 index 00000000..0b1bd2a8 --- /dev/null +++ b/src/routines/level1/xdotu.cc @@ -0,0 +1,49 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren +// +// This file implements the Xdotu class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xdotu.h" + +#include +#include + +namespace clblast { +// ================================================================================================= + +// Constructor: forwards to base class constructor +template +Xdotu::Xdotu(Queue &queue, Event &event, const std::string &name): + Xdot(queue, event, name) { +} + +// ================================================================================================= + +// The main routine +template +StatusCode Xdotu::DoDotu(const size_t n, + const Buffer &dot_buffer, const size_t dot_offset, + const Buffer &x_buffer, const size_t x_offset, const size_t x_inc, + const Buffer &y_buffer, const size_t y_offset, const size_t y_inc) { + return DoDot(n, dot_buffer, dot_offset, + x_buffer, x_offset, x_inc, + y_buffer, y_offset, y_inc, + false); +} + +// ================================================================================================= + +// Compiles the templated class +template class Xdotu; +template class Xdotu; + +// ================================================================================================= +} // namespace clblast diff --git a/src/tuning/xdot.cc b/src/tuning/xdot.cc new file mode 100644 index 00000000..ff6bee16 --- /dev/null +++ b/src/tuning/xdot.cc @@ -0,0 +1,125 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren +// +// This file uses the CLTune auto-tuner to tune the xdot OpenCL kernels. Note that the results are +// not verified, since the result is not final and depends on the WGS2 parameter. +// +// ================================================================================================= + +#include +#include + +#include "internal/utilities.h" +#include "internal/tuning.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template +class TuneXdot { + public: + + // The representative kernel and the source code + static std::string KernelFamily() { return "xdot"; } + static std::string KernelName() { return "Xdot"; } + static std::string GetSources() { + return + #include "../src/kernels/common.opencl" + #include "../src/kernels/level1/xdot.opencl" + ; + } + + // The list of arguments relevant for this routine + static std::vector GetOptions() { return {kArgN}; } + + // Tests for valid arguments + static void TestValidArguments(const Arguments &) { } + + // Sets the default values for the arguments + static size_t DefaultM() { return 1; } // N/A for this kernel + static size_t DefaultN() { return 4096*1024; } + static size_t DefaultK() { return 1; } // N/A for this kernel + static double DefaultFraction() { return 1.0; } // N/A for this kernel + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeX(const Arguments &args) { return args.n; } + static size_t GetSizeY(const Arguments &args) { return args.n; } + static size_t GetSizeA(const Arguments &) { return 1; } // N/A for this kernel + static size_t GetSizeB(const Arguments &) { return 1; } // N/A for this kernel + static size_t GetSizeC(const Arguments &) { return 1; } // N/A for this kernel + static size_t GetSizeTemp(const Arguments &args) { return args.n; } // Worst case + + // Sets the tuning parameters and their possible values + static void SetParameters(cltune::Tuner &tuner, const size_t id) { + tuner.AddParameter(id, "WGS1", {32, 64, 128, 256, 512, 1024}); + tuner.AddParameter(id, "WGS2", {32, 64, 128, 256, 512, 1024}); + tuner.AddParameter(id, "VW", {1}); + } + + // Sets the constraints and local memory size + static void SetConstraints(cltune::Tuner &, const size_t) { } + static void SetLocalMemorySize(cltune::Tuner &, const size_t, const Arguments &) { } + + // Sets the base thread configuration + static std::vector GlobalSize(const Arguments &) { return {2}; } + static std::vector GlobalSizeRef(const Arguments &) { return {2*64*64}; } + static std::vector LocalSize() { return {1}; } + static std::vector LocalSizeRef() { return {64}; } + + // Transforms the thread configuration based on the parameters + using TransformVector = std::vector>; + static TransformVector MulLocal() { return {{"WGS1"}}; } + static TransformVector DivLocal() { return {}; } + static TransformVector MulGlobal() { return {{"WGS1"},{"WGS2"}}; } + static TransformVector DivGlobal() { return {}; } + + // Sets the kernel's arguments + static void SetArguments(cltune::Tuner &tuner, const Arguments &args, + std::vector &x_vec, std::vector &y_vec, + std::vector &, std::vector &, std::vector &, + std::vector &temp) { + tuner.AddArgumentScalar(static_cast(args.n)); + tuner.AddArgumentInput(x_vec); + tuner.AddArgumentScalar(0); + tuner.AddArgumentScalar(1); + tuner.AddArgumentInput(y_vec); + tuner.AddArgumentScalar(0); + tuner.AddArgumentScalar(1); + tuner.AddArgumentInput(temp); // No output checking for the result - size varies + tuner.AddArgumentScalar(static_cast(false)); + } + + // Describes how to compute the performance metrics + static size_t GetMetric(const Arguments &args) { + return (2*args.n + 1) * GetBytes(args.precision); + } + static std::string PerformanceUnit() { return "GB/s"; } +}; + +// ================================================================================================= +} // namespace clblast + +// Shortcuts to the clblast namespace +using float2 = clblast::float2; +using double2 = clblast::double2; + +// Main function (not within the clblast namespace) +int main(int argc, char *argv[]) { + switch(clblast::GetPrecision(argc, argv)) { + case clblast::Precision::kHalf: throw std::runtime_error("Unsupported precision mode"); + case clblast::Precision::kSingle: clblast::Tuner, float>(argc, argv); break; + case clblast::Precision::kDouble: clblast::Tuner, double>(argc, argv); break; + case clblast::Precision::kComplexSingle: clblast::Tuner, float2>(argc, argv); break; + case clblast::Precision::kComplexDouble: clblast::Tuner, double2>(argc, argv); break; + } + return 0; +} + +// ================================================================================================= -- cgit v1.2.3