[arrayfire] 284/408: Adding atomics.hpp file for CUDA that can be used in the future

[Ghislain Vaillant]

This is an automated email from the git hooks/post-receive script.

ghisvail-guest pushed a commit to branch debian/sid
in repository arrayfire.

commit 1fab92387bcb196c607ed4fd1cde19e8a762e1f1
Author: Pavan Yalamanchili <pavan at arrayfire.com>
Date:   Wed Aug 19 01:51:56 2015 -0400

    Adding atomics.hpp file for CUDA that can be used in the future
---
 src/backend/cuda/kernel/atomics.hpp | 59 +++++++++++++++++++++++++++++++++++++
 1 file changed, 59 insertions(+)

diff --git a/src/backend/cuda/kernel/atomics.hpp b/src/backend/cuda/kernel/atomics.hpp
new file mode 100644
index 0000000..1b9ff27
--- /dev/null
+++ b/src/backend/cuda/kernel/atomics.hpp
@@ -0,0 +1,59 @@
+/*******************************************************
+ * Copyright (c) 2015, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+namespace cuda
+{
+    namespace kernel
+    {
+        template<typename T>
+        __device__ T atomicAdd(T *ptr, T val)
+        {
+            return ::atomicAdd(ptr, val);
+        }
+
+#define SPECIALIZE(T, fn1, fn2)                                         \
+        template<>                                                      \
+        __device__ T atomicAdd<T>(T* ptr, T val)                        \
+        {                                                               \
+            unsigned long long int* ptr_as_ull =                        \
+                (unsigned long long int*)ptr;                           \
+            unsigned long long int old = *ptr_as_ull, assumed;          \
+            do {                                                        \
+                assumed = old;                                          \
+                old = atomicCAS(ptr_as_ull, assumed,                    \
+                                fn2(val + fn1(assumed)));               \
+            } while (assumed != old);                                   \
+            return fn1(old);                                            \
+        }                                                               \
+
+        SPECIALIZE(double, __longlong_as_double, __double_as_longlong)
+        SPECIALIZE(intl, intl, uintl)
+        SPECIALIZE(uintl, uintl, uintl)
+
+        template<>
+        __device__ cfloat atomicAdd<cfloat>(cfloat *ptr, cfloat val)
+        {
+            float *fptr = (float *)(ptr);
+            cfloat res;
+            res.x = ::atomicAdd(fptr + 0, val.x);
+            res.y = ::atomicAdd(fptr + 1, val.y);
+            return res;
+        }
+
+        template<>
+        __device__ cdouble atomicAdd<cdouble>(cdouble *ptr, cdouble val)
+        {
+            double *fptr = (double *)(ptr);
+            cdouble res;
+            res.x = atomicAdd(fptr + 0, val.x);
+            res.y = atomicAdd(fptr + 1, val.y);
+            return res;
+        }
+    }
+}

-- 
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/arrayfire.git