From 3876096c30ad4eed5769dbc88dbfe75b7571718a Mon Sep 17 00:00:00 2001 From: Cedric Nugteren Date: Fri, 25 Mar 2016 10:00:40 +0100 Subject: Added prototypes for SNRM2/DNRM2 routines --- src/clblast.cc | 17 +++++++++++++++++ src/clblast_c.cc | 22 ++++++++++++++++++++++ 2 files changed, 39 insertions(+) (limited to 'src') diff --git a/src/clblast.cc b/src/clblast.cc index e7f2477f..9079355a 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -285,6 +285,23 @@ template StatusCode PUBLIC_API Dotc(const size_t, const cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); +// Euclidian norm of a vector: SNRM2/DNRM2 +template +StatusCode Nrm2(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*) { + return StatusCode::kNotImplemented; +} +template StatusCode PUBLIC_API Nrm2(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); +template StatusCode PUBLIC_API Nrm2(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); + // ================================================================================================= // BLAS level-2 (matrix-vector) routines // ================================================================================================= diff --git a/src/clblast_c.cc b/src/clblast_c.cc index 66d16f6d..d735fa35 100644 --- a/src/clblast_c.cc +++ b/src/clblast_c.cc @@ -279,6 +279,28 @@ StatusCode CLBlastZdotc(const size_t n, return static_cast(status); } +// NRM2 +StatusCode CLBlastSnrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Nrm2(n, + nrm2_buffer, nrm2_offset, + x_buffer, x_offset, x_inc, + queue, event); + return static_cast(status); +} +StatusCode CLBlastDnrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Nrm2(n, + nrm2_buffer, nrm2_offset, + x_buffer, x_offset, x_inc, + queue, event); + return static_cast(status); +} + // ================================================================================================= // BLAS level-2 (matrix-vector) routines // ================================================================================================= -- cgit v1.2.3 From 1d5a702d9d31afa320a15ed9fa79471aec314f4a Mon Sep 17 00:00:00 2001 From: Cedric Nugteren Date: Fri, 25 Mar 2016 10:30:38 +0100 Subject: Added prototypes for ScNRM2/DzNRM2 routines --- include/clblast.h | 2 +- include/clblast_c.h | 10 +++++++++- scripts/generator/datatype.py | 3 ++- scripts/generator/generator.py | 28 ++++++++++++++------------- src/clblast.cc | 10 +++++++++- src/clblast_c.cc | 20 +++++++++++++++++++ test/correctness/routines/level1/xnrm2.cc | 2 ++ test/performance/routines/level1/xnrm2.cc | 6 ++++-- test/wrapper_clblas.h | 32 ++++++++++++++++++++++++++++++- 9 files changed, 93 insertions(+), 20 deletions(-) (limited to 'src') diff --git a/include/clblast.h b/include/clblast.h index 905de774..d837cb71 100644 --- a/include/clblast.h +++ b/include/clblast.h @@ -140,7 +140,7 @@ StatusCode Dotc(const size_t n, const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, cl_command_queue* queue, cl_event* event); -// Euclidian norm of a vector: SNRM2/DNRM2 +// Euclidian norm of a vector: SNRM2/DNRM2/ScNRM2/DzNRM2 template StatusCode Nrm2(const size_t n, cl_mem nrm2_buffer, const size_t nrm2_offset, diff --git a/include/clblast_c.h b/include/clblast_c.h index 1e4be1ab..e93ee465 100644 --- a/include/clblast_c.h +++ b/include/clblast_c.h @@ -208,7 +208,7 @@ StatusCode PUBLIC_API CLBlastZdotc(const size_t n, const cl_mem y_buffer, const size_t y_offset, const size_t y_inc, cl_command_queue* queue, cl_event* event); -// Euclidian norm of a vector: SNRM2/DNRM2 +// Euclidian norm of a vector: SNRM2/DNRM2/ScNRM2/DzNRM2 StatusCode PUBLIC_API CLBlastSnrm2(const size_t n, cl_mem nrm2_buffer, const size_t nrm2_offset, const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, @@ -217,6 +217,14 @@ StatusCode PUBLIC_API CLBlastDnrm2(const size_t n, cl_mem nrm2_buffer, const size_t nrm2_offset, const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, cl_command_queue* queue, cl_event* event); +StatusCode PUBLIC_API CLBlastScnrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event); +StatusCode PUBLIC_API CLBlastDznrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event); // ================================================================================================= // BLAS level-2 (matrix-vector) routines diff --git a/scripts/generator/datatype.py b/scripts/generator/datatype.py index 0aa27197..9323bc4d 100644 --- a/scripts/generator/datatype.py +++ b/scripts/generator/datatype.py @@ -22,7 +22,8 @@ D2CL = "cl_double2" # Structure holding data-type and precision information class DataType(): - def __init__(self, name, template, scalars, buffertype): + def __init__(self, precision_name, name, template, scalars, buffertype): + self.precision_name = precision_name self.name = name self.template = template self.alpha_cpp = scalars[0] diff --git a/scripts/generator/generator.py b/scripts/generator/generator.py index 2c22a6fd..253f1a92 100644 --- a/scripts/generator/generator.py +++ b/scripts/generator/generator.py @@ -31,21 +31,23 @@ from datatype import DataType, FLT, DBL, FLT2, DBL2, F2CL, D2CL # ================================================================================================== # Regular data-types -S = DataType("S", FLT, [FLT, FLT, FLT, FLT], FLT ) # single (32) -D = DataType("D", DBL, [DBL, DBL, DBL, DBL], DBL ) # double (64) -C = DataType("C", FLT2, [FLT2, FLT2, F2CL, F2CL], FLT2) # single-complex (3232) -Z = DataType("Z", DBL2, [DBL2, DBL2, D2CL, D2CL], DBL2) # double-complex (6464) +S = DataType("S", "S", FLT, [FLT, FLT, FLT, FLT], FLT ) # single (32) +D = DataType("D", "D", DBL, [DBL, DBL, DBL, DBL], DBL ) # double (64) +C = DataType("C", "C", FLT2, [FLT2, FLT2, F2CL, F2CL], FLT2) # single-complex (3232) +Z = DataType("Z", "Z", DBL2, [DBL2, DBL2, D2CL, D2CL], DBL2) # double-complex (6464) # Special cases -Css = DataType("C", FLT, [FLT, FLT, FLT, FLT], FLT2) # As C, but with constants from S -Zdd = DataType("Z", DBL, [DBL, DBL, DBL, DBL], DBL2) # As Z, but with constants from D -Ccs = DataType("C", FLT2+","+FLT, [FLT2, FLT, F2CL, FLT], FLT2) # As C, but with one constant from S -Zzd = DataType("Z", DBL2+","+DBL, [DBL2, DBL, D2CL, DBL], DBL2) # As Z, but with one constant from D +Sc = DataType("C", "Sc", FLT2, [FLT2, FLT2, FLT2, FLT2], FLT2) # As C, but with real output +Dz = DataType("Z", "Dz", DBL2, [DBL2, DBL2, DBL2, DBL2], DBL2) # As Z, but with real output +Css = DataType("C", "C", FLT, [FLT, FLT, FLT, FLT], FLT2) # As C, but with constants from S +Zdd = DataType("Z", "Z", DBL, [DBL, DBL, DBL, DBL], DBL2) # As Z, but with constants from D +Ccs = DataType("C", "C", FLT2+","+FLT, [FLT2, FLT, F2CL, FLT], FLT2) # As C, but with one constant from S +Zzd = DataType("Z", "Z", DBL2+","+DBL, [DBL2, DBL, D2CL, DBL], DBL2) # As Z, but with one constant from D # C++ template data-types -T = DataType("typename T", "T", ["T", "T", "T", "T"], "T") # regular routine -Tc = DataType("typename T", "std::complex,T", ["T", "T", "T", "T"], "std::complex") # for herk -TU = DataType("typename T, typename U", "T,U", ["T", "U", "T", "U"], "T") # for her2k +T = DataType("T", "typename T", "T", ["T", "T", "T", "T"], "T") # regular routine +Tc = DataType("Tc", "typename T", "std::complex,T", ["T", "T", "T", "T"], "std::complex") # for herk +TU = DataType("TU", "typename T, typename U", "T,U", ["T", "U", "T", "U"], "T") # for her2k # ================================================================================================== @@ -61,7 +63,7 @@ routines = [ Routine(True, "1", "dot", T, [S,D], ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two vectors"), Routine(True, "1", "dotu", T, [C,Z], ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two complex vectors"), Routine(True, "1", "dotc", T, [C,Z], ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two complex vectors, one conjugated"), - Routine(False, "1", "nrm2", T, [S,D], ["n"], [], ["x"], ["nrm2"], [], True, "Euclidian norm of a vector"), + Routine(False, "1", "nrm2", T, [S,D,Sc,Dz],["n"], [], ["x"], ["nrm2"], [], True, "Euclidian norm of a vector"), ], [ # Level 2: matrix-vector Routine(True, "2a", "gemv", T, [S,D,C,Z], ["m","n"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], False, "General matrix-vector multiplication"), @@ -332,7 +334,7 @@ for level in [1,2,3]: body += " case clblast::Precision::k"+PrecisionToFullName(precision)+":" found = False for flavour in routine.flavours: - if flavour.name == precision: + if flavour.precision_name == precision: body += "\n clblast::RunClient(const size_t, const cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); -// Euclidian norm of a vector: SNRM2/DNRM2 +// Euclidian norm of a vector: SNRM2/DNRM2/ScNRM2/DzNRM2 template StatusCode Nrm2(const size_t, cl_mem, const size_t, @@ -301,6 +301,14 @@ template StatusCode PUBLIC_API Nrm2(const size_t, cl_mem, const size_t, const cl_mem, const size_t, const size_t, cl_command_queue*, cl_event*); +template StatusCode PUBLIC_API Nrm2(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); +template StatusCode PUBLIC_API Nrm2(const size_t, + cl_mem, const size_t, + const cl_mem, const size_t, const size_t, + cl_command_queue*, cl_event*); // ================================================================================================= // BLAS level-2 (matrix-vector) routines diff --git a/src/clblast_c.cc b/src/clblast_c.cc index d735fa35..fa25d4a7 100644 --- a/src/clblast_c.cc +++ b/src/clblast_c.cc @@ -300,6 +300,26 @@ StatusCode CLBlastDnrm2(const size_t n, queue, event); return static_cast(status); } +StatusCode CLBlastScnrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Nrm2(n, + nrm2_buffer, nrm2_offset, + x_buffer, x_offset, x_inc, + queue, event); + return static_cast(status); +} +StatusCode CLBlastDznrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto status = clblast::Nrm2(n, + nrm2_buffer, nrm2_offset, + x_buffer, x_offset, x_inc, + queue, event); + return static_cast(status); +} // ================================================================================================= // BLAS level-2 (matrix-vector) routines diff --git a/test/correctness/routines/level1/xnrm2.cc b/test/correctness/routines/level1/xnrm2.cc index 8238e868..97fb0ad6 100644 --- a/test/correctness/routines/level1/xnrm2.cc +++ b/test/correctness/routines/level1/xnrm2.cc @@ -20,6 +20,8 @@ using double2 = clblast::double2; int main(int argc, char *argv[]) { clblast::RunTests, float, float>(argc, argv, false, "SNRM2"); clblast::RunTests, double, double>(argc, argv, true, "DNRM2"); + clblast::RunTests, float2, float2>(argc, argv, true, "ScNRM2"); + clblast::RunTests, double2, double2>(argc, argv, true, "DzNRM2"); return 0; } diff --git a/test/performance/routines/level1/xnrm2.cc b/test/performance/routines/level1/xnrm2.cc index d5ae348b..db6ec9ad 100644 --- a/test/performance/routines/level1/xnrm2.cc +++ b/test/performance/routines/level1/xnrm2.cc @@ -24,8 +24,10 @@ int main(int argc, char *argv[]) { clblast::RunClient, float, float>(argc, argv); break; case clblast::Precision::kDouble: clblast::RunClient, double, double>(argc, argv); break; - case clblast::Precision::kComplexSingle: throw std::runtime_error("Unsupported precision mode"); - case clblast::Precision::kComplexDouble: throw std::runtime_error("Unsupported precision mode"); + 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/wrapper_clblas.h b/test/wrapper_clblas.h index 501f0bc5..37d9eee5 100644 --- a/test/wrapper_clblas.h +++ b/test/wrapper_clblas.h @@ -350,7 +350,7 @@ clblasStatus clblasXdotc(const size_t n, num_queues, queues, num_wait_events, wait_events, events); } -// Forwards the clBLAS calls for SNRM2/DNRM2 +// Forwards the clBLAS calls for SNRM2/DNRM2/ScNRM2/DzNRM2 template clblasStatus clblasXnrm2(const size_t n, cl_mem nrm2_buffer, const size_t nrm2_offset, @@ -387,6 +387,36 @@ clblasStatus clblasXnrm2(const size_t n, scratch_buffer(), num_queues, queues, num_wait_events, wait_events, events); } +template <> +clblasStatus clblasXnrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto queue = Queue(queues[0]); + auto context = queue.GetContext(); + auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + return clblasScnrm2(n, + nrm2_buffer, nrm2_offset, + x_buffer, x_offset, static_cast(x_inc), + scratch_buffer(), + num_queues, queues, num_wait_events, wait_events, events); +} +template <> +clblasStatus clblasXnrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_uint num_queues, cl_command_queue *queues, + cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { + auto queue = Queue(queues[0]); + auto context = queue.GetContext(); + auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + return clblasDznrm2(n, + nrm2_buffer, nrm2_offset, + x_buffer, x_offset, static_cast(x_inc), + scratch_buffer(), + num_queues, queues, num_wait_events, wait_events, events); +} // ================================================================================================= // BLAS level-2 (matrix-vector) routines -- cgit v1.2.3 From aaa687ca984b18bd1ea499c92285b490fd78e2a3 Mon Sep 17 00:00:00 2001 From: Cedric Nugteren Date: Mon, 28 Mar 2016 23:00:44 +0200 Subject: Added preliminary support for the xNRM2 routines --- CMakeLists.txt | 2 +- include/internal/routines/level1/xnrm2.h | 55 ++++++++++++++ include/internal/utilities.h | 6 +- scripts/generator/generator.py | 92 ++++++++++++------------ src/clblast.cc | 18 +++-- src/kernels/level1/xnrm2.opencl | 120 +++++++++++++++++++++++++++++++ src/routines/level1/xnrm2.cc | 107 +++++++++++++++++++++++++++ test/correctness/testblas.cc | 28 ++++---- test/correctness/testblas.h | 14 ++-- test/performance/client.cc | 11 +-- test/routines/level1/xdot.h | 10 +-- test/routines/level1/xdotc.h | 10 +-- test/routines/level1/xdotu.h | 10 +-- test/routines/level1/xnrm2.h | 117 ++++++++++++++++++++++++++++++ test/wrapper_clblas.h | 44 ++++++------ 15 files changed, 529 insertions(+), 115 deletions(-) create mode 100644 include/internal/routines/level1/xnrm2.h create mode 100644 src/kernels/level1/xnrm2.opencl create mode 100644 src/routines/level1/xnrm2.cc create mode 100644 test/routines/level1/xnrm2.h (limited to 'src') diff --git a/CMakeLists.txt b/CMakeLists.txt index 65debdf4..db73c83e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -118,7 +118,7 @@ include_directories(${clblast_SOURCE_DIR}/include ${OPENCL_INCLUDE_DIRS}) set(KERNELS copy pad transpose padtranspose xaxpy xdot xger xgemm xgemv) set(SAMPLE_PROGRAMS_CPP sgemm) set(SAMPLE_PROGRAMS_C sgemm) -set(LEVEL1_ROUTINES xswap xscal xcopy xaxpy xdot xdotu xdotc) +set(LEVEL1_ROUTINES xswap xscal xcopy xaxpy xdot xdotu xdotc xnrm2) set(LEVEL2_ROUTINES xgemv xgbmv xhemv xhbmv xhpmv xsymv xsbmv xspmv xtrmv xtbmv xtpmv xger xgeru xgerc xher xhpr xher2 xhpr2 xsyr xspr xsyr2 xspr2) set(LEVEL3_ROUTINES xgemm xsymm xhemm xsyrk xherk xsyr2k xher2k xtrmm) diff --git a/include/internal/routines/level1/xnrm2.h b/include/internal/routines/level1/xnrm2.h new file mode 100644 index 00000000..b3fffef6 --- /dev/null +++ b/include/internal/routines/level1/xnrm2.h @@ -0,0 +1,55 @@ + +// ================================================================================================= +// 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 Xnrm2 routine. The precision is implemented using a template argument. +// +// ================================================================================================= + +#ifndef CLBLAST_ROUTINES_XNRM2_H_ +#define CLBLAST_ROUTINES_XNRM2_H_ + +#include "internal/routine.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template +class Xnrm2: public Routine { + public: + + // Members and methods from the base class + using Routine::db_; + using Routine::source_string_; + using Routine::queue_; + using Routine::context_; + using Routine::GetProgramFromCache; + using Routine::TestVectorX; + using Routine::TestVectorDot; + using Routine::RunKernel; + using Routine::ErrorIn; + + // Constructor + Xnrm2(Queue &queue, Event &event, const std::string &name = "NRM2"); + + // Templated-precision implementation of the routine + StatusCode DoNrm2(const size_t n, + const Buffer &nrm2_buffer, const size_t nrm2_offset, + const Buffer &x_buffer, const size_t x_offset, const size_t x_inc); + + private: + // Static variable to get the precision + const static Precision precision_; +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_ROUTINES_XNRM2_H_ +#endif diff --git a/include/internal/utilities.h b/include/internal/utilities.h index b6307a85..35f76722 100644 --- a/include/internal/utilities.h +++ b/include/internal/utilities.h @@ -61,6 +61,7 @@ constexpr auto kArgBOffset = "offb"; constexpr auto kArgCOffset = "offc"; constexpr auto kArgAPOffset = "offap"; constexpr auto kArgDotOffset = "offdot"; +constexpr auto kArgNrm2Offset = "offnrm2"; constexpr auto kArgAlpha = "alpha"; constexpr auto kArgBeta = "beta"; @@ -113,6 +114,7 @@ struct Arguments { size_t c_offset = 0; size_t ap_offset = 0; size_t dot_offset = 0; + size_t nrm2_offset = 0; T alpha = T{1.0}; T beta = T{1.0}; size_t x_size = 1; @@ -121,7 +123,7 @@ struct Arguments { size_t b_size = 1; size_t c_size = 1; size_t ap_size = 1; - size_t dot_size = 1; + size_t scalar_size = 1; // Tuner-specific arguments double fraction = 1.0; // Client-specific arguments @@ -149,7 +151,7 @@ struct Buffers { Buffer b_mat; Buffer c_mat; Buffer ap_mat; - Buffer dot; + Buffer scalar; }; // ================================================================================================= diff --git a/scripts/generator/generator.py b/scripts/generator/generator.py index 253f1a92..c316bea6 100644 --- a/scripts/generator/generator.py +++ b/scripts/generator/generator.py @@ -54,55 +54,55 @@ TU = DataType("TU", "typename T, typename U", "T,U", ["T", "U", "T", "U"], "T") # Populates a list of routines routines = [ [ # Level 1: vector-vector - #Routine(False, "1", "rotg", T, [S,D], [], [], [], [], ["a","b","c","s"], False, "Generate plane rotation"), - #Routine(False, "1", "rot", T, [S,D], ["n"], [], [], ["x","y"], ["c","s"], False, "Apply plane rotation"), - Routine(True, "1", "swap", T, [S,D,C,Z], ["n"], [], [], ["x","y"], [], False, "Swap two vectors"), - Routine(True, "1", "scal", T, [S,D,C,Z], ["n"], [], [], ["x"], ["alpha"], False, "Vector scaling"), - Routine(True, "1", "copy", T, [S,D,C,Z], ["n"], [], ["x"], ["y"], [], False, "Vector copy"), - Routine(True, "1", "axpy", T, [S,D,C,Z], ["n"], [], ["x"], ["y"], ["alpha"], False, "Vector-times-constant plus vector"), - Routine(True, "1", "dot", T, [S,D], ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two vectors"), - Routine(True, "1", "dotu", T, [C,Z], ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two complex vectors"), - Routine(True, "1", "dotc", T, [C,Z], ["n"], [], ["x","y"], ["dot"], [], True, "Dot product of two complex vectors, one conjugated"), - Routine(False, "1", "nrm2", T, [S,D,Sc,Dz],["n"], [], ["x"], ["nrm2"], [], True, "Euclidian norm of a vector"), + #Routine(False, "1", "rotg", T, [S,D], [], [], [], [], ["a","b","c","s"], "", "Generate plane rotation"), + #Routine(False, "1", "rot", T, [S,D], ["n"], [], [], ["x","y"], ["c","s"], "", "Apply plane rotation"), + Routine(True, "1", "swap", T, [S,D,C,Z], ["n"], [], [], ["x","y"], [], "", "Swap two vectors"), + Routine(True, "1", "scal", T, [S,D,C,Z], ["n"], [], [], ["x"], ["alpha"], "", "Vector scaling"), + Routine(True, "1", "copy", T, [S,D,C,Z], ["n"], [], ["x"], ["y"], [], "", "Vector copy"), + Routine(True, "1", "axpy", T, [S,D,C,Z], ["n"], [], ["x"], ["y"], ["alpha"], "", "Vector-times-constant plus vector"), + Routine(True, "1", "dot", T, [S,D], ["n"], [], ["x","y"], ["dot"], [], "n", "Dot product of two vectors"), + Routine(True, "1", "dotu", T, [C,Z], ["n"], [], ["x","y"], ["dot"], [], "n", "Dot product of two complex vectors"), + Routine(True, "1", "dotc", T, [C,Z], ["n"], [], ["x","y"], ["dot"], [], "n", "Dot product of two complex vectors, one conjugated"), + Routine(True, "1", "nrm2", T, [S,D,Sc,Dz],["n"], [], ["x"], ["nrm2"], [], "2*n", "Euclidian norm of a vector"), ], [ # Level 2: matrix-vector - Routine(True, "2a", "gemv", T, [S,D,C,Z], ["m","n"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], False, "General matrix-vector multiplication"), - Routine(True, "2a", "gbmv", T, [S,D,C,Z], ["m","n","kl","ku"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], False, "General banded matrix-vector multiplication"), - Routine(True, "2a", "hemv", T, [C,Z], ["n"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], False, "Hermitian matrix-vector multiplication"), - Routine(True, "2a", "hbmv", T, [C,Z], ["n","k"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], False, "Hermitian banded matrix-vector multiplication"), - Routine(True, "2a", "hpmv", T, [C,Z], ["n"], ["layout","triangle"], ["ap","x"], ["y"], ["alpha","beta"], False, "Hermitian packed matrix-vector multiplication"), - Routine(True, "2a", "symv", T, [S,D], ["n"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], False, "Symmetric matrix-vector multiplication"), - Routine(True, "2a", "sbmv", T, [S,D], ["n","k"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], False, "Symmetric banded matrix-vector multiplication"), - Routine(True, "2a", "spmv", T, [S,D], ["n"], ["layout","triangle"], ["ap","x"], ["y"], ["alpha","beta"], False, "Symmetric packed matrix-vector multiplication"), - Routine(True, "2a", "trmv", T, [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], True, "Triangular matrix-vector multiplication"), - Routine(True, "2a", "tbmv", T, [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], True, "Triangular banded matrix-vector multiplication"), - Routine(True, "2a", "tpmv", T, [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["ap"], ["x"], [], True, "Triangular packed matrix-vector multiplication"), - Routine(False, "2a", "trsv", T, [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], False, "Solves a triangular system of equations"), - Routine(False, "2a", "tbsv", T, [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], False, "Solves a banded triangular system of equations"), - Routine(False, "2a", "tpsv", T, [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["ap"], ["x"], [], False, "Solves a packed triangular system of equations"), + Routine(True, "2a", "gemv", T, [S,D,C,Z], ["m","n"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], "", "General matrix-vector multiplication"), + Routine(True, "2a", "gbmv", T, [S,D,C,Z], ["m","n","kl","ku"], ["layout","a_transpose"], ["a","x"], ["y"], ["alpha","beta"], "", "General banded matrix-vector multiplication"), + Routine(True, "2a", "hemv", T, [C,Z], ["n"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], "", "Hermitian matrix-vector multiplication"), + Routine(True, "2a", "hbmv", T, [C,Z], ["n","k"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], "", "Hermitian banded matrix-vector multiplication"), + Routine(True, "2a", "hpmv", T, [C,Z], ["n"], ["layout","triangle"], ["ap","x"], ["y"], ["alpha","beta"], "", "Hermitian packed matrix-vector multiplication"), + Routine(True, "2a", "symv", T, [S,D], ["n"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], "", "Symmetric matrix-vector multiplication"), + Routine(True, "2a", "sbmv", T, [S,D], ["n","k"], ["layout","triangle"], ["a","x"], ["y"], ["alpha","beta"], "", "Symmetric banded matrix-vector multiplication"), + Routine(True, "2a", "spmv", T, [S,D], ["n"], ["layout","triangle"], ["ap","x"], ["y"], ["alpha","beta"], "", "Symmetric packed matrix-vector multiplication"), + Routine(True, "2a", "trmv", T, [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], "n", "Triangular matrix-vector multiplication"), + Routine(True, "2a", "tbmv", T, [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], "n", "Triangular banded matrix-vector multiplication"), + Routine(True, "2a", "tpmv", T, [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["ap"], ["x"], [], "n", "Triangular packed matrix-vector multiplication"), + Routine(False, "2a", "trsv", T, [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], "", "Solves a triangular system of equations"), + Routine(False, "2a", "tbsv", T, [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose","diagonal"], ["a"], ["x"], [], "", "Solves a banded triangular system of equations"), + Routine(False, "2a", "tpsv", T, [S,D,C,Z], ["n"], ["layout","triangle","a_transpose","diagonal"], ["ap"], ["x"], [], "", "Solves a packed triangular system of equations"), # Level 2: matrix update - Routine(True, "2b", "ger", T, [S,D], ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 matrix update"), - Routine(True, "2b", "geru", T, [C,Z], ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 complex matrix update"), - Routine(True, "2b", "gerc", T, [C,Z], ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], False, "General rank-1 complex conjugated matrix update"), - Routine(True, "2b", "her", Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], False, "Hermitian rank-1 matrix update"), - Routine(True, "2b", "hpr", Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], False, "Hermitian packed rank-1 matrix update"), - Routine(True, "2b", "her2", T, [C,Z], ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], False, "Hermitian rank-2 matrix update"), - Routine(True, "2b", "hpr2", T, [C,Z], ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], False, "Hermitian packed rank-2 matrix update"), - Routine(True, "2b", "syr", T, [S,D], ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], False, "Symmetric rank-1 matrix update"), - Routine(True, "2b", "spr", T, [S,D], ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], False, "Symmetric packed rank-1 matrix update"), - Routine(True, "2b", "syr2", T, [S,D], ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], False, "Symmetric rank-2 matrix update"), - Routine(True, "2b", "spr2", T, [S,D], ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], False, "Symmetric packed rank-2 matrix update"), + Routine(True, "2b", "ger", T, [S,D], ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], "", "General rank-1 matrix update"), + Routine(True, "2b", "geru", T, [C,Z], ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], "", "General rank-1 complex matrix update"), + Routine(True, "2b", "gerc", T, [C,Z], ["m","n"], ["layout"], ["x","y"], ["a"], ["alpha"], "", "General rank-1 complex conjugated matrix update"), + Routine(True, "2b", "her", Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], "", "Hermitian rank-1 matrix update"), + Routine(True, "2b", "hpr", Tc, [Css,Zdd], ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], "", "Hermitian packed rank-1 matrix update"), + Routine(True, "2b", "her2", T, [C,Z], ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], "", "Hermitian rank-2 matrix update"), + Routine(True, "2b", "hpr2", T, [C,Z], ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], "", "Hermitian packed rank-2 matrix update"), + Routine(True, "2b", "syr", T, [S,D], ["n"], ["layout","triangle"], ["x"], ["a"], ["alpha"], "", "Symmetric rank-1 matrix update"), + Routine(True, "2b", "spr", T, [S,D], ["n"], ["layout","triangle"], ["x"], ["ap"], ["alpha"], "", "Symmetric packed rank-1 matrix update"), + Routine(True, "2b", "syr2", T, [S,D], ["n"], ["layout","triangle"], ["x","y"], ["a"], ["alpha"], "", "Symmetric rank-2 matrix update"), + Routine(True, "2b", "spr2", T, [S,D], ["n"], ["layout","triangle"], ["x","y"], ["ap"], ["alpha"], "", "Symmetric packed rank-2 matrix update"), ], [ # Level 3: matrix-matrix - Routine(True, "3", "gemm", T, [S,D,C,Z], ["m","n","k"], ["layout","a_transpose","b_transpose"], ["a","b"], ["c"], ["alpha","beta"], False, "General matrix-matrix multiplication"), - Routine(True, "3", "symm", T, [S,D,C,Z], ["m","n"], ["layout","side","triangle"], ["a","b"], ["c"], ["alpha","beta"], False, "Symmetric matrix-matrix multiplication"), - Routine(True, "3", "hemm", T, [C,Z], ["m","n"], ["layout","side","triangle"], ["a","b"], ["c"], ["alpha","beta"], False, "Hermitian matrix-matrix multiplication"), - Routine(True, "3", "syrk", T, [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose"], ["a"], ["c"], ["alpha","beta"], False, "Rank-K update of a symmetric matrix"), - Routine(True, "3", "herk", Tc, [Css,Zdd], ["n","k"], ["layout","triangle","a_transpose"], ["a"], ["c"], ["alpha","beta"], False, "Rank-K update of a hermitian matrix"), - Routine(True, "3", "syr2k", T, [S,D,C,Z], ["n","k"], ["layout","triangle","ab_transpose"], ["a","b"], ["c"], ["alpha","beta"], False, "Rank-2K update of a symmetric matrix"), - Routine(True, "3", "her2k", TU, [Ccs,Zzd], ["n","k"], ["layout","triangle","ab_transpose"], ["a","b"], ["c"], ["alpha","beta"], False, "Rank-2K update of a hermitian matrix"), - Routine(True, "3", "trmm", T, [S,D,C,Z], ["m","n"], ["layout","side","triangle","a_transpose","diagonal"], ["a"], ["b"], ["alpha"], False, "Triangular matrix-matrix multiplication"), - Routine(False, "3", "trsm", T, [S,D,C,Z], ["m","n"], ["layout","side","triangle","a_transpose","diagonal"], ["a"], ["b"], ["alpha"], False, "Solves a triangular system of equations"), + Routine(True, "3", "gemm", T, [S,D,C,Z], ["m","n","k"], ["layout","a_transpose","b_transpose"], ["a","b"], ["c"], ["alpha","beta"], "", "General matrix-matrix multiplication"), + Routine(True, "3", "symm", T, [S,D,C,Z], ["m","n"], ["layout","side","triangle"], ["a","b"], ["c"], ["alpha","beta"], "", "Symmetric matrix-matrix multiplication"), + Routine(True, "3", "hemm", T, [C,Z], ["m","n"], ["layout","side","triangle"], ["a","b"], ["c"], ["alpha","beta"], "", "Hermitian matrix-matrix multiplication"), + Routine(True, "3", "syrk", T, [S,D,C,Z], ["n","k"], ["layout","triangle","a_transpose"], ["a"], ["c"], ["alpha","beta"], "", "Rank-K update of a symmetric matrix"), + Routine(True, "3", "herk", Tc, [Css,Zdd], ["n","k"], ["layout","triangle","a_transpose"], ["a"], ["c"], ["alpha","beta"], "", "Rank-K update of a hermitian matrix"), + Routine(True, "3", "syr2k", T, [S,D,C,Z], ["n","k"], ["layout","triangle","ab_transpose"], ["a","b"], ["c"], ["alpha","beta"], "", "Rank-2K update of a symmetric matrix"), + Routine(True, "3", "her2k", TU, [Ccs,Zzd], ["n","k"], ["layout","triangle","ab_transpose"], ["a","b"], ["c"], ["alpha","beta"], "", "Rank-2K update of a hermitian matrix"), + Routine(True, "3", "trmm", T, [S,D,C,Z], ["m","n"], ["layout","side","triangle","a_transpose","diagonal"], ["a"], ["b"], ["alpha"], "", "Triangular matrix-matrix multiplication"), + Routine(False, "3", "trsm", T, [S,D,C,Z], ["m","n"], ["layout","side","triangle","a_transpose","diagonal"], ["a"], ["b"], ["alpha"], "", "Solves a triangular system of equations"), ]] # ================================================================================================== @@ -226,7 +226,7 @@ def wrapper_clblas(routines): if routine.scratch: result += " auto queue = Queue(queues[0]);\n" result += " auto context = queue.GetContext();\n" - result += " auto scratch_buffer = Buffer<"+flavour.template+">(context, n*x_inc + x_offset);\n" + result += " auto scratch_buffer = Buffer<"+flavour.template+">(context, "+routine.scratch+");\n" arguments += ["scratch_buffer()"] result += " return clblas"+flavour.name+routine.name+"(" result += (",\n"+indent).join([a for a in arguments]) @@ -250,7 +250,7 @@ files = [ path_clblast+"/src/clblast_c.cc", path_clblast+"/test/wrapper_clblas.h", ] -header_lines = [84, 64, 93, 22, 22] +header_lines = [84, 65, 93, 22, 22] footer_lines = [6, 3, 9, 2, 6] # Checks whether the command-line arguments are valid; exists otherwise diff --git a/src/clblast.cc b/src/clblast.cc index 2c940380..6f75540d 100644 --- a/src/clblast.cc +++ b/src/clblast.cc @@ -26,6 +26,7 @@ #include "internal/routines/level1/xdot.h" #include "internal/routines/level1/xdotu.h" #include "internal/routines/level1/xdotc.h" +#include "internal/routines/level1/xnrm2.h" // BLAS level-2 includes #include "internal/routines/level2/xgemv.h" @@ -287,11 +288,18 @@ template StatusCode PUBLIC_API Dotc(const size_t, // Euclidian norm of a vector: SNRM2/DNRM2/ScNRM2/DzNRM2 template -StatusCode Nrm2(const size_t, - cl_mem, const size_t, - const cl_mem, const size_t, const size_t, - cl_command_queue*, cl_event*) { - return StatusCode::kNotImplemented; +StatusCode Nrm2(const size_t n, + cl_mem nrm2_buffer, const size_t nrm2_offset, + const cl_mem x_buffer, const size_t x_offset, const size_t x_inc, + cl_command_queue* queue, cl_event* event) { + auto queue_cpp = Queue(*queue); + auto event_cpp = Event(*event); + auto routine = Xnrm2(queue_cpp, event_cpp); + auto status = routine.SetUp(); + if (status != StatusCode::kSuccess) { return status; } + return routine.DoNrm2(n, + Buffer(nrm2_buffer), nrm2_offset, + Buffer(x_buffer), x_offset, x_inc); } template StatusCode PUBLIC_API Nrm2(const size_t, cl_mem, const size_t, diff --git a/src/kernels/level1/xnrm2.opencl b/src/kernels/level1/xnrm2.opencl new file mode 100644 index 00000000..c50d7d63 --- /dev/null +++ b/src/kernels/level1/xnrm2.opencl @@ -0,0 +1,120 @@ + +// ================================================================================================= +// 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 Xnrm2 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 Xnrm2(const int n, + const __global real* restrict xgm, const int x_offset, const int x_inc, + __global real* output) { + __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 x1 = xgm[id*x_inc + x_offset]; + real x2 = x1; + COMPLEX_CONJUGATE(x2); + MultiplyAdd(acc, x1, x2); + 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]; + } +} + +// ================================================================================================= + +// Computes the square root +inline real SquareRoot(const real z) { + #if PRECISION == 3232 || PRECISION == 6464 + double r = sqrt(z.x * z.x + z.y * z.y); + real zpr; zpr.x = z.x + r; zpr.y = z.y; + double zprabs = sqrt(zpr.x * zpr.x + zpr.y + zpr.y); + real result; + result.x = sqrt(r) * zpr.x / zprabs; + result.y = sqrt(r) * zpr.y / zprabs; + return result; + #else + return sqrt(z); + #endif +} + +// 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 Xnrm2Epilogue(const __global real* restrict input, + __global real* nrm2, const int nrm2_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); + } + + // Computes the square root and stores the final result + if (lid == 0) { + nrm2[nrm2_offset] = SquareRoot(lm[0]); + } +} + +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= diff --git a/src/routines/level1/xnrm2.cc b/src/routines/level1/xnrm2.cc new file mode 100644 index 00000000..064e68bf --- /dev/null +++ b/src/routines/level1/xnrm2.cc @@ -0,0 +1,107 @@ + +// ================================================================================================= +// 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 Xnrm2 class (see the header for information about the class). +// +// ================================================================================================= + +#include "internal/routines/level1/xnrm2.h" + +#include +#include + +namespace clblast { +// ================================================================================================= + +// Specific implementations to get the memory-type based on a template argument +template <> const Precision Xnrm2::precision_ = Precision::kSingle; +template <> const Precision Xnrm2::precision_ = Precision::kDouble; +template <> const Precision Xnrm2::precision_ = Precision::kComplexSingle; +template <> const Precision Xnrm2::precision_ = Precision::kComplexDouble; + +// ================================================================================================= + +// Constructor: forwards to base class constructor +template +Xnrm2::Xnrm2(Queue &queue, Event &event, const std::string &name): + Routine(queue, event, name, {"Xdot"}, precision_) { + source_string_ = + #include "../../kernels/level1/xnrm2.opencl" + ; +} + +// ================================================================================================= + +// The main routine +template +StatusCode Xnrm2::DoNrm2(const size_t n, + const Buffer &nrm2_buffer, const size_t nrm2_offset, + const Buffer &x_buffer, const size_t x_offset, const size_t x_inc) { + + // Makes sure all dimensions are larger than zero + if (n == 0) { return StatusCode::kInvalidDimension; } + + // Tests the vectors for validity + auto status = TestVectorX(n, x_buffer, x_offset, x_inc, sizeof(T)); + if (ErrorIn(status)) { return status; } + status = TestVectorDot(1, nrm2_buffer, nrm2_offset, sizeof(T)); + if (ErrorIn(status)) { return status; } + + // Retrieves the Xnrm2 kernels from the compiled binary + try { + auto& program = GetProgramFromCache(); + auto kernel1 = Kernel(program, "Xnrm2"); + auto kernel2 = Kernel(program, "Xnrm2Epilogue"); + + // 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, temp_buffer()); + + // 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, nrm2_buffer()); + kernel2.SetArgument(2, static_cast(nrm2_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 Xnrm2; +template class Xnrm2; +template class Xnrm2; +template class Xnrm2; + +// ================================================================================================= +} // namespace clblast diff --git a/test/correctness/testblas.cc b/test/correctness/testblas.cc index febd7504..1329b2c5 100644 --- a/test/correctness/testblas.cc +++ b/test/correctness/testblas.cc @@ -58,14 +58,14 @@ TestBlas::TestBlas(int argc, char *argv[], const bool silent, b_source_.resize(std::max(max_mat, max_matvec)*std::max(max_ld, max_matvec) + max_offset); c_source_.resize(std::max(max_mat, max_matvec)*std::max(max_ld, max_matvec) + max_offset); ap_source_.resize(std::max(max_mat, max_matvec)*std::max(max_mat, max_matvec) + max_offset); - dot_source_.resize(std::max(max_mat, max_matvec) + max_offset); + scalar_source_.resize(std::max(max_mat, max_matvec) + max_offset); PopulateVector(x_source_); PopulateVector(y_source_); PopulateVector(a_source_); PopulateVector(b_source_); PopulateVector(c_source_); PopulateVector(ap_source_); - PopulateVector(dot_source_); + PopulateVector(scalar_source_); } // =============================================================================================== @@ -86,15 +86,15 @@ void TestBlas::TestRegular(std::vector> &test_vector, const st auto b_mat1 = Buffer(context_, args.b_size); auto c_mat1 = Buffer(context_, args.c_size); auto ap_mat1 = Buffer(context_, args.ap_size); - auto dot1 = Buffer(context_, args.dot_size); + auto scalar1 = Buffer(context_, args.scalar_size); x_vec1.Write(queue_, args.x_size, x_source_); y_vec1.Write(queue_, args.y_size, y_source_); a_mat1.Write(queue_, args.a_size, a_source_); b_mat1.Write(queue_, args.b_size, b_source_); c_mat1.Write(queue_, args.c_size, c_source_); ap_mat1.Write(queue_, args.ap_size, ap_source_); - dot1.Write(queue_, args.dot_size, dot_source_); - auto buffers1 = Buffers{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, ap_mat1, dot1}; + scalar1.Write(queue_, args.scalar_size, scalar_source_); + auto buffers1 = Buffers{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, ap_mat1, scalar1}; auto status1 = run_reference_(args, buffers1, queue_); // Runs the CLBlast code @@ -104,15 +104,15 @@ void TestBlas::TestRegular(std::vector> &test_vector, const st auto b_mat2 = Buffer(context_, args.b_size); auto c_mat2 = Buffer(context_, args.c_size); auto ap_mat2 = Buffer(context_, args.ap_size); - auto dot2 = Buffer(context_, args.dot_size); + auto scalar2 = Buffer(context_, args.scalar_size); x_vec2.Write(queue_, args.x_size, x_source_); y_vec2.Write(queue_, args.y_size, y_source_); a_mat2.Write(queue_, args.a_size, a_source_); b_mat2.Write(queue_, args.b_size, b_source_); c_mat2.Write(queue_, args.c_size, c_source_); ap_mat2.Write(queue_, args.ap_size, ap_source_); - dot2.Write(queue_, args.dot_size, dot_source_); - auto buffers2 = Buffers{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, ap_mat2, dot2}; + scalar2.Write(queue_, args.scalar_size, scalar_source_); + auto buffers2 = Buffers{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, ap_mat2, scalar2}; auto status2 = run_routine_(args, buffers2, queue_); // Tests for equality of the two status codes @@ -162,32 +162,32 @@ void TestBlas::TestInvalid(std::vector> &test_vector, const st auto b1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.b_size*sizeof(T), nullptr,nullptr); auto c1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.c_size*sizeof(T), nullptr,nullptr); auto ap1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.ap_size*sizeof(T), nullptr,nullptr); - auto d1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.dot_size*sizeof(T), nullptr,nullptr); + auto d1 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.scalar_size*sizeof(T), nullptr,nullptr); auto x_vec1 = Buffer(x1); auto y_vec1 = Buffer(y1); auto a_mat1 = Buffer(a1); auto b_mat1 = Buffer(b1); auto c_mat1 = Buffer(c1); auto ap_mat1 = Buffer(ap1); - auto dot1 = Buffer(d1); + auto scalar1 = Buffer(d1); auto x2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.x_size*sizeof(T), nullptr,nullptr); auto y2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.y_size*sizeof(T), nullptr,nullptr); auto a2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.a_size*sizeof(T), nullptr,nullptr); auto b2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.b_size*sizeof(T), nullptr,nullptr); auto c2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.c_size*sizeof(T), nullptr,nullptr); auto ap2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.ap_size*sizeof(T), nullptr,nullptr); - auto d2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.dot_size*sizeof(T), nullptr,nullptr); + auto d2 = clCreateBuffer(context_(), CL_MEM_READ_WRITE, args.scalar_size*sizeof(T), nullptr,nullptr); auto x_vec2 = Buffer(x2); auto y_vec2 = Buffer(y2); auto a_mat2 = Buffer(a2); auto b_mat2 = Buffer(b2); auto c_mat2 = Buffer(c2); auto ap_mat2 = Buffer(ap2); - auto dot2 = Buffer(d2); + auto scalar2 = Buffer(d2); // Runs the two routines - auto buffers1 = Buffers{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, ap_mat1, dot1}; - auto buffers2 = Buffers{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, ap_mat2, dot2}; + auto buffers1 = Buffers{x_vec1, y_vec1, a_mat1, b_mat1, c_mat1, ap_mat1, scalar1}; + auto buffers2 = Buffers{x_vec2, y_vec2, a_mat2, b_mat2, c_mat2, ap_mat2, scalar2}; auto status1 = run_reference_(args, buffers1, queue_); auto status2 = run_routine_(args, buffers2, queue_); diff --git a/test/correctness/testblas.h b/test/correctness/testblas.h index bfd1763c..7c9032bd 100644 --- a/test/correctness/testblas.h +++ b/test/correctness/testblas.h @@ -92,7 +92,7 @@ class TestBlas: public Tester { std::vector b_source_; std::vector c_source_; std::vector ap_source_; - std::vector dot_source_; + std::vector scalar_source_; // The routine-specific functions passed to the tester Routine run_routine_; @@ -143,6 +143,7 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name auto c_offsets = std::vector{args.c_offset}; auto ap_offsets = std::vector{args.ap_offset}; auto dot_offsets = std::vector{args.dot_offset}; + auto nrm2_offsets = std::vector{args.nrm2_offset}; auto alphas = std::vector{args.alpha}; auto betas = std::vector{args.beta}; auto x_sizes = std::vector{args.x_size}; @@ -182,6 +183,7 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name if (option == kArgCOffset) { c_offsets = tester.kOffsets; } if (option == kArgAPOffset) { ap_offsets = tester.kOffsets; } if (option == kArgDotOffset) { dot_offsets = tester.kOffsets; } + if (option == kArgNrm2Offset) { nrm2_offsets = tester.kOffsets; } if (option == kArgAlpha) { alphas = tester.kAlphaValues; } if (option == kArgBeta) { betas = tester.kBetaValues; } @@ -221,10 +223,12 @@ void RunTests(int argc, char *argv[], const bool silent, const std::string &name for (auto &c_offset: c_offsets) { r_args.c_offset = c_offset; for (auto &ap_offset: ap_offsets) { r_args.ap_offset = ap_offset; for (auto &dot_offset: dot_offsets) { r_args.dot_offset = dot_offset; - for (auto &alpha: alphas) { r_args.alpha = alpha; - for (auto &beta: betas) { r_args.beta = beta; - C::SetSizes(r_args); - regular_test_vector.push_back(r_args); + for (auto &nrm2_offset: nrm2_offsets) { r_args.nrm2_offset = nrm2_offset; + for (auto &alpha: alphas) { r_args.alpha = alpha; + for (auto &beta: betas) { r_args.beta = beta; + C::SetSizes(r_args); + regular_test_vector.push_back(r_args); + } } } } diff --git a/test/performance/client.cc b/test/performance/client.cc index ebfad3a6..17f54231 100644 --- a/test/performance/client.cc +++ b/test/performance/client.cc @@ -136,14 +136,14 @@ void Client::PerformanceTest(Arguments &args, const SetMetric set_sizes) std::vector b_source(args.b_size); std::vector c_source(args.c_size); std::vector ap_source(args.ap_size); - std::vector dot_source(args.dot_size); + std::vector scalar_source(args.scalar_size); PopulateVector(x_source); PopulateVector(y_source); PopulateVector(a_source); PopulateVector(b_source); PopulateVector(c_source); PopulateVector(ap_source); - PopulateVector(dot_source); + PopulateVector(scalar_source); // Creates the matrices on the device auto x_vec = Buffer(context, args.x_size); @@ -152,15 +152,15 @@ void Client::PerformanceTest(Arguments &args, const SetMetric set_sizes) auto b_mat = Buffer(context, args.b_size); auto c_mat = Buffer(context, args.c_size); auto ap_mat = Buffer(context, args.ap_size); - auto dot = Buffer(context, args.dot_size); + auto scalar = Buffer(context, args.scalar_size); x_vec.Write(queue, args.x_size, x_source); y_vec.Write(queue, args.y_size, y_source); a_mat.Write(queue, args.a_size, a_source); b_mat.Write(queue, args.b_size, b_source); c_mat.Write(queue, args.c_size, c_source); ap_mat.Write(queue, args.ap_size, ap_source); - dot.Write(queue, args.dot_size, dot_source); - auto buffers = Buffers{x_vec, y_vec, a_mat, b_mat, c_mat, ap_mat, dot}; + scalar.Write(queue, args.scalar_size, scalar_source); + auto buffers = Buffers{x_vec, y_vec, a_mat, b_mat, c_mat, ap_mat, scalar}; // Runs the routines and collects the timings auto timings = std::vector>(); @@ -267,6 +267,7 @@ void Client::PrintTableRow(const Arguments& args, else if (o == kArgCOffset) { integers.push_back(args.c_offset); } else if (o == kArgAPOffset) { integers.push_back(args.ap_offset); } else if (o == kArgDotOffset) {integers.push_back(args.dot_offset); } + else if (o == kArgNrm2Offset){integers.push_back(args.nrm2_offset); } } auto strings = std::vector{}; for (auto &o: options_) { diff --git a/test/routines/level1/xdot.h b/test/routines/level1/xdot.h index bfcfdaff..04669f52 100644 --- a/test/routines/level1/xdot.h +++ b/test/routines/level1/xdot.h @@ -54,7 +54,7 @@ class TestXdot { static void SetSizes(Arguments &args) { args.x_size = GetSizeX(args); args.y_size = GetSizeY(args); - args.dot_size = GetSizeDot(args); + args.scalar_size = GetSizeDot(args); } // Describes what the default values of the leading dimensions of the matrices are @@ -72,7 +72,7 @@ class TestXdot { auto queue_plain = queue(); auto event = cl_event{}; auto status = Dot(args.n, - buffers.dot(), args.dot_offset, + buffers.scalar(), args.dot_offset, buffers.x_vec(), args.x_offset, args.x_inc, buffers.y_vec(), args.y_offset, args.y_inc, &queue_plain, &event); @@ -85,7 +85,7 @@ class TestXdot { auto queue_plain = queue(); auto event = cl_event{}; auto status = clblasXdot(args.n, - buffers.dot(), args.dot_offset, + buffers.scalar(), args.dot_offset, buffers.x_vec(), args.x_offset, args.x_inc, buffers.y_vec(), args.y_offset, args.y_inc, 1, &queue_plain, 0, nullptr, &event); @@ -95,8 +95,8 @@ class TestXdot { // Describes how to download the results of the computation (more importantly: which buffer) static std::vector DownloadResult(const Arguments &args, Buffers &buffers, Queue &queue) { - std::vector result(args.dot_size, static_cast(0)); - buffers.dot.Read(queue, args.dot_size, result); + std::vector result(args.scalar_size, static_cast(0)); + buffers.scalar.Read(queue, args.scalar_size, result); return result; } diff --git a/test/routines/level1/xdotc.h b/test/routines/level1/xdotc.h index e403ba4c..e5b42ef4 100644 --- a/test/routines/level1/xdotc.h +++ b/test/routines/level1/xdotc.h @@ -54,7 +54,7 @@ class TestXdotc { static void SetSizes(Arguments &args) { args.x_size = GetSizeX(args); args.y_size = GetSizeY(args); - args.dot_size = GetSizeDot(args); + args.scalar_size = GetSizeDot(args); } // Describes what the default values of the leading dimensions of the matrices are @@ -72,7 +72,7 @@ class TestXdotc { auto queue_plain = queue(); auto event = cl_event{}; auto status = Dotc(args.n, - buffers.dot(), args.dot_offset, + buffers.scalar(), args.dot_offset, buffers.x_vec(), args.x_offset, args.x_inc, buffers.y_vec(), args.y_offset, args.y_inc, &queue_plain, &event); @@ -85,7 +85,7 @@ class TestXdotc { auto queue_plain = queue(); auto event = cl_event{}; auto status = clblasXdotc(args.n, - buffers.dot(), args.dot_offset, + buffers.scalar(), args.dot_offset, buffers.x_vec(), args.x_offset, args.x_inc, buffers.y_vec(), args.y_offset, args.y_inc, 1, &queue_plain, 0, nullptr, &event); @@ -95,8 +95,8 @@ class TestXdotc { // Describes how to download the results of the computation (more importantly: which buffer) static std::vector DownloadResult(const Arguments &args, Buffers &buffers, Queue &queue) { - std::vector result(args.dot_size, static_cast(0)); - buffers.dot.Read(queue, args.dot_size, result); + std::vector result(args.scalar_size, static_cast(0)); + buffers.scalar.Read(queue, args.scalar_size, result); return result; } diff --git a/test/routines/level1/xdotu.h b/test/routines/level1/xdotu.h index 8b2c65a8..6430148c 100644 --- a/test/routines/level1/xdotu.h +++ b/test/routines/level1/xdotu.h @@ -54,7 +54,7 @@ class TestXdotu { static void SetSizes(Arguments &args) { args.x_size = GetSizeX(args); args.y_size = GetSizeY(args); - args.dot_size = GetSizeDot(args); + args.scalar_size = GetSizeDot(args); } // Describes what the default values of the leading dimensions of the matrices are @@ -72,7 +72,7 @@ class TestXdotu { auto queue_plain = queue(); auto event = cl_event{}; auto status = Dotu(args.n, - buffers.dot(), args.dot_offset, + buffers.scalar(), args.dot_offset, buffers.x_vec(), args.x_offset, args.x_inc, buffers.y_vec(), args.y_offset, args.y_inc, &queue_plain, &event); @@ -85,7 +85,7 @@ class TestXdotu { auto queue_plain = queue(); auto event = cl_event{}; auto status = clblasXdotu(args.n, - buffers.dot(), args.dot_offset, + buffers.scalar(), args.dot_offset, buffers.x_vec(), args.x_offset, args.x_inc, buffers.y_vec(), args.y_offset, args.y_inc, 1, &queue_plain, 0, nullptr, &event); @@ -95,8 +95,8 @@ class TestXdotu { // Describes how to download the results of the computation (more importantly: which buffer) static std::vector DownloadResult(const Arguments &args, Buffers &buffers, Queue &queue) { - std::vector result(args.dot_size, static_cast(0)); - buffers.dot.Read(queue, args.dot_size, result); + std::vector result(args.scalar_size, static_cast(0)); + buffers.scalar.Read(queue, args.scalar_size, result); return result; } diff --git a/test/routines/level1/xnrm2.h b/test/routines/level1/xnrm2.h new file mode 100644 index 00000000..e3f77ee4 --- /dev/null +++ b/test/routines/level1/xnrm2.h @@ -0,0 +1,117 @@ + +// ================================================================================================= +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Cedric Nugteren +// +// This file implements a class with static methods to describe the Xnrm2 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_XNRM2_H_ +#define CLBLAST_TEST_ROUTINES_XNRM2_H_ + +#include +#include + +#include "wrapper_clblas.h" + +namespace clblast { +// ================================================================================================= + +// See comment at top of file for a description of the class +template +class TestXnrm2 { + public: + + // The BLAS level: 1, 2, or 3 + static size_t BLASLevel() { return 1; } + + // The list of arguments relevant for this routine + static std::vector GetOptions() { + return {kArgN, + kArgXInc, + kArgXOffset, kArgNrm2Offset}; + } + + // Describes how to obtain the sizes of the buffers + static size_t GetSizeX(const Arguments &args) { + return args.n * args.x_inc + args.x_offset; + } + static size_t GetSizeNrm2(const Arguments &args) { + return 1 + args.nrm2_offset; + } + + // Describes how to set the sizes of all the buffers + static void SetSizes(Arguments &args) { + args.x_size = GetSizeX(args); + args.scalar_size = GetSizeNrm2(args); + } + + // Describes what the default values of the leading dimensions of the matrices are + static size_t DefaultLDA(const Arguments &) { return 1; } // N/A for this routine + static size_t DefaultLDB(const Arguments &) { return 1; } // N/A for this routine + static size_t DefaultLDC(const Arguments &) { return 1; } // N/A for this routine + + // Describes which transpose options are relevant for this routine + using Transposes = std::vector; + static Transposes GetATransposes(const Transposes &) { return {}; } // N/A for this routine + static Transposes GetBTransposes(const Transposes &) { return {}; } // N/A for this routine + + // Describes how to run the CLBlast routine + static StatusCode RunRoutine(const Arguments &args, const Buffers &buffers, Queue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = Nrm2(args.n, + buffers.scalar(), args.nrm2_offset, + buffers.x_vec(), args.x_offset, args.x_inc, + &queue_plain, &event); + clWaitForEvents(1, &event); + return status; + } + + // Describes how to run the clBLAS routine (for correctness/performance comparison) + static StatusCode RunReference(const Arguments &args, const Buffers &buffers, Queue &queue) { + auto queue_plain = queue(); + auto event = cl_event{}; + auto status = clblasXnrm2(args.n, + buffers.scalar(), args.nrm2_offset, + buffers.x_vec(), args.x_offset, args.x_inc, + 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, Queue &queue) { + std::vector result(args.scalar_size, static_cast(0)); + buffers.scalar.Read(queue, args.scalar_size, result); + return result; + } + + // Describes how to compute the indices of the result buffer + static size_t ResultID1(const Arguments &) { return 1; } // N/A for this routine + static size_t ResultID2(const Arguments &) { return 1; } // N/A for this routine + static size_t GetResultIndex(const Arguments &args, const size_t, const size_t) { + return args.nrm2_offset; + } + + // Describes how to compute performance metrics + static size_t GetFlops(const Arguments &args) { + return 2 * args.n; + } + static size_t GetBytes(const Arguments &args) { + return ((args.n) + 1) * sizeof(T); + } +}; + +// ================================================================================================= +} // namespace clblast + +// CLBLAST_TEST_ROUTINES_XNRM2_H_ +#endif diff --git a/test/wrapper_clblas.h b/test/wrapper_clblas.h index 37d9eee5..42bb8f92 100644 --- a/test/wrapper_clblas.h +++ b/test/wrapper_clblas.h @@ -238,7 +238,7 @@ clblasStatus clblasXdot(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasSdot(n, dot_buffer, dot_offset, x_buffer, x_offset, static_cast(x_inc), @@ -255,7 +255,7 @@ clblasStatus clblasXdot(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasDdot(n, dot_buffer, dot_offset, x_buffer, x_offset, static_cast(x_inc), @@ -281,7 +281,7 @@ clblasStatus clblasXdotu(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasCdotu(n, dot_buffer, dot_offset, x_buffer, x_offset, static_cast(x_inc), @@ -298,7 +298,7 @@ clblasStatus clblasXdotu(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasZdotu(n, dot_buffer, dot_offset, x_buffer, x_offset, static_cast(x_inc), @@ -324,7 +324,7 @@ clblasStatus clblasXdotc(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasCdotc(n, dot_buffer, dot_offset, x_buffer, x_offset, static_cast(x_inc), @@ -341,7 +341,7 @@ clblasStatus clblasXdotc(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasZdotc(n, dot_buffer, dot_offset, x_buffer, x_offset, static_cast(x_inc), @@ -365,7 +365,7 @@ clblasStatus clblasXnrm2(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, 2*n); return clblasSnrm2(n, nrm2_buffer, nrm2_offset, x_buffer, x_offset, static_cast(x_inc), @@ -380,7 +380,7 @@ clblasStatus clblasXnrm2(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, 2*n); return clblasDnrm2(n, nrm2_buffer, nrm2_offset, x_buffer, x_offset, static_cast(x_inc), @@ -395,7 +395,7 @@ clblasStatus clblasXnrm2(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, 2*n); return clblasScnrm2(n, nrm2_buffer, nrm2_offset, x_buffer, x_offset, static_cast(x_inc), @@ -410,7 +410,7 @@ clblasStatus clblasXnrm2(const size_t n, cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, 2*n); return clblasDznrm2(n, nrm2_buffer, nrm2_offset, x_buffer, x_offset, static_cast(x_inc), @@ -815,7 +815,7 @@ clblasStatus clblasXtrmv(const clblasOrder layout, const clblasUplo trian cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasStrmv(layout, triangle, a_transpose, diagonal, n, a_buffer, a_offset, a_ld, @@ -832,7 +832,7 @@ clblasStatus clblasXtrmv(const clblasOrder layout, const clblasUplo tria cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasDtrmv(layout, triangle, a_transpose, diagonal, n, a_buffer, a_offset, a_ld, @@ -849,7 +849,7 @@ clblasStatus clblasXtrmv(const clblasOrder layout, const clblasUplo tria cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasCtrmv(layout, triangle, a_transpose, diagonal, n, a_buffer, a_offset, a_ld, @@ -866,7 +866,7 @@ clblasStatus clblasXtrmv(const clblasOrder layout, const clblasUplo tri cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasZtrmv(layout, triangle, a_transpose, diagonal, n, a_buffer, a_offset, a_ld, @@ -892,7 +892,7 @@ clblasStatus clblasXtbmv(const clblasOrder layout, const clblasUplo trian cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasStbmv(layout, triangle, a_transpose, diagonal, n, k, a_buffer, a_offset, a_ld, @@ -909,7 +909,7 @@ clblasStatus clblasXtbmv(const clblasOrder layout, const clblasUplo tria cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasDtbmv(layout, triangle, a_transpose, diagonal, n, k, a_buffer, a_offset, a_ld, @@ -926,7 +926,7 @@ clblasStatus clblasXtbmv(const clblasOrder layout, const clblasUplo tria cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasCtbmv(layout, triangle, a_transpose, diagonal, n, k, a_buffer, a_offset, a_ld, @@ -943,7 +943,7 @@ clblasStatus clblasXtbmv(const clblasOrder layout, const clblasUplo tri cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasZtbmv(layout, triangle, a_transpose, diagonal, n, k, a_buffer, a_offset, a_ld, @@ -969,7 +969,7 @@ clblasStatus clblasXtpmv(const clblasOrder layout, const clblasUplo trian cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasStpmv(layout, triangle, a_transpose, diagonal, n, ap_buffer, ap_offset, @@ -986,7 +986,7 @@ clblasStatus clblasXtpmv(const clblasOrder layout, const clblasUplo tria cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasDtpmv(layout, triangle, a_transpose, diagonal, n, ap_buffer, ap_offset, @@ -1003,7 +1003,7 @@ clblasStatus clblasXtpmv(const clblasOrder layout, const clblasUplo tria cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasCtpmv(layout, triangle, a_transpose, diagonal, n, ap_buffer, ap_offset, @@ -1020,7 +1020,7 @@ clblasStatus clblasXtpmv(const clblasOrder layout, const clblasUplo tri cl_uint num_wait_events, const cl_event *wait_events, cl_event *events) { auto queue = Queue(queues[0]); auto context = queue.GetContext(); - auto scratch_buffer = Buffer(context, n*x_inc + x_offset); + auto scratch_buffer = Buffer(context, n); return clblasZtpmv(layout, triangle, a_transpose, diagonal, n, ap_buffer, ap_offset, -- cgit v1.2.3 From 5409f349a17f60ba68133fd0cc9789fb2918f790 Mon Sep 17 00:00:00 2001 From: cnugteren Date: Wed, 30 Mar 2016 21:32:04 -0700 Subject: Fixed the nrm2 kernel for complex data-types --- src/kernels/level1/xnrm2.opencl | 21 +++++---------------- 1 file changed, 5 insertions(+), 16 deletions(-) (limited to 'src') diff --git a/src/kernels/level1/xnrm2.opencl b/src/kernels/level1/xnrm2.opencl index c50d7d63..cf579457 100644 --- a/src/kernels/level1/xnrm2.opencl +++ b/src/kernels/level1/xnrm2.opencl @@ -70,21 +70,6 @@ __kernel void Xnrm2(const int n, // ================================================================================================= -// Computes the square root -inline real SquareRoot(const real z) { - #if PRECISION == 3232 || PRECISION == 6464 - double r = sqrt(z.x * z.x + z.y * z.y); - real zpr; zpr.x = z.x + r; zpr.y = z.y; - double zprabs = sqrt(zpr.x * zpr.x + zpr.y + zpr.y); - real result; - result.x = sqrt(r) * zpr.x / zprabs; - result.y = sqrt(r) * zpr.y / zprabs; - return result; - #else - return sqrt(z); - #endif -} - // 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))) @@ -108,7 +93,11 @@ __kernel void Xnrm2Epilogue(const __global real* restrict input, // Computes the square root and stores the final result if (lid == 0) { - nrm2[nrm2_offset] = SquareRoot(lm[0]); + #if PRECISION == 3232 || PRECISION == 6464 + nrm2[nrm2_offset].x = sqrt(lm[0].x); // the result is a non-complex number + #else + nrm2[nrm2_offset] = sqrt(lm[0]); + #endif } } -- cgit v1.2.3