diff options
author | CNugteren <web@cedricnugteren.nl> | 2015-09-14 16:57:00 +0200 |
---|---|---|
committer | CNugteren <web@cedricnugteren.nl> | 2015-09-14 16:57:00 +0200 |
commit | a2e726d3bd4294f1eae1735f6ba23105dccc6b10 (patch) | |
tree | 2685c3204460a300088b13af66128cdda8e64c68 /src/kernels | |
parent | 2a383f34501b386b8e6c4beb56c6ac694622f060 (diff) |
Added xDOT/xDOTU/xDOTC dot-product routines
Diffstat (limited to 'src/kernels')
-rw-r--r-- | src/kernels/common.opencl | 7 | ||||
-rw-r--r-- | src/kernels/level1/xdot.opencl | 106 |
2 files changed, 113 insertions, 0 deletions
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 <www.cedricnugteren.nl> +// +// 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 +)" + +// ================================================================================================= |