diff options
author | cnugteren <web@cedricnugteren.nl> | 2016-04-20 22:12:51 -0600 |
---|---|---|
committer | cnugteren <web@cedricnugteren.nl> | 2016-04-20 22:12:51 -0600 |
commit | 16a048f1ac8102ad4bcce9cf9fc320f791323e45 (patch) | |
tree | bd7df539c92b92cb8841ff27fd83c4e14e5958bc /src/kernels | |
parent | 894983fc3c7c57ffc48c21523641694cde318eca (diff) |
Added support for the iSAMAX/iDAMAX/iCAMAX/iZAMAX routines
Diffstat (limited to 'src/kernels')
-rw-r--r-- | src/kernels/common.opencl | 9 | ||||
-rw-r--r-- | src/kernels/level1/xamax.opencl | 128 |
2 files changed, 137 insertions, 0 deletions
diff --git a/src/kernels/common.opencl b/src/kernels/common.opencl index 0a68defb..57d75ee0 100644 --- a/src/kernels/common.opencl +++ b/src/kernels/common.opencl @@ -80,6 +80,15 @@ R"( #define ONE 1.0 #endif +// Single-element version of a complex number +#if PRECISION == 3232 + typedef float singlereal; +#elif PRECISION == 6464 + typedef double singlereal; +#else + typedef real singlereal; +#endif + // ================================================================================================= // Don't use the non-IEEE754 compliant OpenCL built-in mad() instruction per default. For specific diff --git a/src/kernels/level1/xamax.opencl b/src/kernels/level1/xamax.opencl new file mode 100644 index 00000000..03dd05e5 --- /dev/null +++ b/src/kernels/level1/xamax.opencl @@ -0,0 +1,128 @@ + +// ================================================================================================= +// 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 Xamax kernel. It implements an index of absolute max computation using +// reduction kernels. Reduction is split in two parts. In the first (main) kernel the X vector is +// loaded, 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 loading and the majority of the operation +__attribute__((reqd_work_group_size(WGS1, 1, 1))) +__kernel void Xamax(const int n, + const __global real* restrict xgm, const int x_offset, const int x_inc, + __global singlereal* maxgm, __global unsigned int* imaxgm) { + __local singlereal maxlm[WGS1]; + __local unsigned int imaxlm[WGS1]; + const int lid = get_local_id(0); + const int wgid = get_group_id(0); + const int num_groups = get_num_groups(0); + + // Performs loading and the first steps of the reduction + singlereal max = ZERO; + unsigned int imax = 0; + int id = wgid*WGS1 + lid; + while (id < n) { + #if PRECISION == 3232 || PRECISION == 6464 + singlereal x = fabs(xgm[id*x_inc + x_offset].x); + #else + singlereal x = fabs(xgm[id*x_inc + x_offset]); + #endif + if (x >= max) { + max = x; + imax = id*x_inc + x_offset; + } + id += WGS1*num_groups; + } + maxlm[lid] = max; + imaxlm[lid] = imax; + 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) { + if (maxlm[lid + s] >= maxlm[lid]) { + maxlm[lid] = maxlm[lid + s]; + imaxlm[lid] = imaxlm[lid + s]; + } + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + // Stores the per-workgroup result + if (lid == 0) { + maxgm[wgid] = maxlm[0]; + imaxgm[wgid] = imaxlm[0]; + } +} + +// ================================================================================================= + +// The epilogue reduction kernel, performing the final bit of the operation. This kernel has to +// be launched with a single workgroup only. +__attribute__((reqd_work_group_size(WGS2, 1, 1))) +__kernel void XamaxEpilogue(const __global singlereal* restrict maxgm, + const __global unsigned int* restrict imaxgm, + __global unsigned int* imax, const int imax_offset) { + __local singlereal maxlm[WGS2]; + __local unsigned int imaxlm[WGS2]; + const int lid = get_local_id(0); + + // Performs the first step of the reduction while loading the data + if (maxgm[lid + WGS2] >= maxgm[lid]) { + maxlm[lid] = maxgm[lid + WGS2]; + imaxlm[lid] = imaxgm[lid + WGS2]; + } + else { + maxlm[lid] = maxgm[lid]; + imaxlm[lid] = imaxgm[lid]; + } + 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) { + if (maxlm[lid + s] >= maxlm[lid]) { + maxlm[lid] = maxlm[lid + s]; + imaxlm[lid] = imaxlm[lid + s]; + } + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + // Stores the final result + if (lid == 0) { + imax[imax_offset] = imaxlm[0]; + } +} + +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= |