From 8be99de82d2ff0634c1289d9b4d1785364a68a44 Mon Sep 17 00:00:00 2001 From: cnugteren Date: Thu, 14 Apr 2016 19:58:26 -0600 Subject: Added support for the SASUM/DASUM/ScASUM/DzASUM routines --- src/kernels/common.opencl | 7 +++ src/kernels/level1/xasum.opencl | 108 ++++++++++++++++++++++++++++++++++++++++ src/kernels/level1/xnrm2.opencl | 10 ++-- 3 files changed, 120 insertions(+), 5 deletions(-) create mode 100644 src/kernels/level1/xasum.opencl (limited to 'src/kernels') diff --git a/src/kernels/common.opencl b/src/kernels/common.opencl index f2a2e7a7..0a68defb 100644 --- a/src/kernels/common.opencl +++ b/src/kernels/common.opencl @@ -109,6 +109,13 @@ R"( #define SetToOne(a) a = ONE #endif +// The absolute value (component-wise) +#if PRECISION == 3232 || PRECISION == 6464 + #define AbsoluteValue(value) value.x = fabs(value.x); value.y = fabs(value.y) +#else + #define AbsoluteValue(value) value = fabs(value) +#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 diff --git a/src/kernels/level1/xasum.opencl b/src/kernels/level1/xasum.opencl new file mode 100644 index 00000000..037dc57e --- /dev/null +++ b/src/kernels/level1/xasum.opencl @@ -0,0 +1,108 @@ + +// ================================================================================================= +// 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 Xasum kernel. It implements a absolute sum 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 Xasum(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 loading 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]; + AbsoluteValue(x); + Add(acc, acc, x); + 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 operation. This kernel has to +// be launched with a single workgroup only. +__attribute__((reqd_work_group_size(WGS2, 1, 1))) +__kernel void XasumEpilogue(const __global real* restrict input, + __global real* asum, const int asum_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 absolute value and stores the final result + if (lid == 0) { + #if PRECISION == 3232 || PRECISION == 6464 + asum[asum_offset].x = lm[0].x + lm[0].y; // the result is a non-complex number + #else + asum[asum_offset] = lm[0]; + #endif + } +} + +// ================================================================================================= + +// End of the C++11 raw string literal +)" + +// ================================================================================================= diff --git a/src/kernels/level1/xnrm2.opencl b/src/kernels/level1/xnrm2.opencl index cf579457..9803687a 100644 --- a/src/kernels/level1/xnrm2.opencl +++ b/src/kernels/level1/xnrm2.opencl @@ -7,9 +7,9 @@ // 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 +// This file contains the Xnrm2 kernel. It implements a squared norm computation using reduction +// kernels. Reduction is split in two parts. In the first (main) kernel the X vector is squared, +// 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. // // ================================================================================================= @@ -29,7 +29,7 @@ R"( // ================================================================================================= -// The main reduction kernel, performing the multiplication and the majority of the sum operation +// The main reduction kernel, performing the multiplication and the majority of the 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, @@ -70,7 +70,7 @@ __kernel void Xnrm2(const int n, // ================================================================================================= -// The epilogue reduction kernel, performing the final bit of the sum operation. This kernel has to +// 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 Xnrm2Epilogue(const __global real* restrict input, -- cgit v1.2.3