diff options
author | cnugteren <web@cedricnugteren.nl> | 2016-03-30 21:32:04 -0700 |
---|---|---|
committer | cnugteren <web@cedricnugteren.nl> | 2016-03-30 21:32:04 -0700 |
commit | 5409f349a17f60ba68133fd0cc9789fb2918f790 (patch) | |
tree | 96a61d9358cd161942d7ccedbf71a98639787eee | |
parent | aaa687ca984b18bd1ea499c92285b490fd78e2a3 (diff) |
Fixed the nrm2 kernel for complex data-types
-rw-r--r-- | src/kernels/level1/xnrm2.opencl | 21 |
1 files changed, 5 insertions, 16 deletions
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 } } |