From 2baac9cd8bc1995a00653036f8c033963f932814 Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Wed, 10 Feb 2016 07:51:31 -0800
Subject: [PATCH 01/27] initial commit of atomic adds

---
 vtkm/cont/DeviceAdapterAlgorithm.h            |  8 ++
 .../internal/DeviceAdapterAlgorithmCuda.h     | 56 +++++++++++++
 .../internal/DeviceAdapterAlgorithmGeneral.h  | 34 ++++++++
 .../tbb/internal/DeviceAdapterAlgorithmTBB.h  | 60 +++++++++++++
 vtkm/cont/testing/TestingDeviceAdapter.h      | 84 +++++++++++++++++++
 vtkm/exec/AtomicArray.h                       | 63 ++++++++++++++
 6 files changed, 305 insertions(+)
 create mode 100644 vtkm/exec/AtomicArray.h

diff --git a/vtkm/cont/DeviceAdapterAlgorithm.h b/vtkm/cont/DeviceAdapterAlgorithm.h
index b8f6e1d61..92528995d 100644
--- a/vtkm/cont/DeviceAdapterAlgorithm.h
+++ b/vtkm/cont/DeviceAdapterAlgorithm.h
@@ -523,6 +523,14 @@ public:
   }
 };
 
+/// \brief Class providing a device-specific support for atomic operations.
+///
+/// The class provide the actual implementation used by
+/// vtkm::cont::DeviceAdapterAtomicArrayImplementation.
+///
+template<typename T, typename DeviceTag>
+class DeviceAdapterAtomicArrayImplementation;
+
 }
 } // namespace vtkm::cont
 
diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
index ca66b27d9..54cd62924 100644
--- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
+++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
@@ -185,6 +185,62 @@ private:
   vtkm::Int32 HighestArchSupported;
 };
 
+/// CUDA contains its own atomic operations
+///
+template<typename T>
+class DeviceAdapterAtomicArrayImplementation<T,vtkm::cont::DeviceAdapterTagCuda>
+{
+public:
+  VTKM_CONT_EXPORT
+  DeviceAdapterAtomicArrayImplementation(
+             vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> handle):
+    Portal( handle.PrepareForInPlace( vtkm::cont::DeviceAdapterTagCuda()) )
+  {
+  }
+
+  inline __device__
+  T Add(vtkm::Id index, const T& value) const
+  {
+    T *lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
+    return vtkmAtomicAdd(lockedValue, value);
+  }
+
+private:
+  typedef typename vtkm::cont::ArrayHandle<T,vtkm::cont::StorageTagBasic>
+        ::template ExecutionTypes<vtkm::cont::DeviceAdapterTagCuda>::Portal PortalType;
+  PortalType Portal;
+
+  inline __device__
+  vtkm::Int64 vtkmAtomicAdd(vtkm::Int64 *address, const vtkm::Int64 &value) const
+  {
+    return atomicAdd((unsigned long long *)address,(unsigned long long) value);
+  }
+
+  inline __device__
+  vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
+  {
+    return atomicAdd((unsigned long long *)address,(unsigned long long) value);
+  }
+  
+  inline __device__
+  vtkm::Float32 vtkmAtomicAdd(vtkm::Float32 *address, const vtkm::Float32 &value) const
+  {
+    return atomicAdd(address,value);
+  }
+
+  inline __device__
+  vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
+  {
+    return atomicAdd(address,value);
+  }
+
+  inline __device__
+  vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
+  {
+    return atomicAdd(address,value);
+  }  
+};
+
 }
 } // namespace vtkm::cont
 
diff --git a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
index 19bb8b940..9ce72a3c6 100644
--- a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
+++ b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
@@ -695,4 +695,38 @@ template<typename T, typename U, class CIn, class CStencil, class COut>
 }
 } // namespace vtkm::cont::internal
 
+namespace vtkm {
+namespace cont {
+/// \brief Class providing a device-specific atomic interface.
+///
+/// The class provide the actual implementation used by vtkm::exec::AtomicArray.
+/// A serial default implementation is provided. But each device will have a different
+/// implementation.
+///
+/// Serial requires no form of atomicity
+///
+template<typename T, typename DeviceTag>
+class DeviceAdapterAtomicArrayImplementation
+{
+public:
+  VTKM_CONT_EXPORT
+  DeviceAdapterAtomicArrayImplementation(vtkm::cont::ArrayHandle<T> handle):
+    Portal( handle.PrepareForInPlace(DeviceTag()) )
+  {
+  }
+
+  VTKM_EXEC_EXPORT
+  T Add(vtkm::Id index, const T& value) const
+  {
+    const T old = this->Portal.Get(index);
+    this->Portal.Set(index, old + value);
+    return old;
+  }
+
+private:
+    typedef typename vtkm::cont::ArrayHandle<T>
+              ::template ExecutionTypes<DeviceTag>::Portal PortalType;
+  PortalType Portal;
+};
+
 #endif //vtk_m_cont_internal_DeviceAdapterAlgorithmGeneral_h
diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index ea833ad0f..7d53d2cff 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -297,6 +297,66 @@ private:
   ::tbb::tick_count StartTime;
 };
 
+template<typename T>
+class DeviceAdapterAtomicArrayImplementation<T,vtkm::cont::DeviceAdapterTagTBB>
+{
+public:
+  VTKM_CONT_EXPORT
+  DeviceAdapterAtomicArrayImplementation(
+               vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> handle):
+    Iterators( IteratorsType( handle.PrepareForInPlace(
+                                      vtkm::cont::DeviceAdapterTagTBB())
+                             ) )
+  {
+  }
+
+  VTKM_EXEC_EXPORT
+  T Add(vtkm::Id index, const T& value) const
+  {
+    T* lockedValue; 
+    lockedValue = (Iterators.GetBegin()+index);
+    return vtkmAtomicAdd(lockedValue, value);
+  }
+
+private:
+  typedef typename vtkm::cont::ArrayHandle<T,vtkm::cont::StorageTagBasic>
+        ::template ExecutionTypes<DeviceAdapterTagTBB>::Portal PortalType;
+  typedef vtkm::cont::ArrayPortalToIterators<PortalType> IteratorsType;
+
+  IteratorsType Iterators;
+
+  VTKM_EXEC_EXPORT
+  vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
+  {
+    return __sync_fetch_and_add(address,value);
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::Int64 vtkmAtomicAdd(vtkm::Int64 *address, const vtkm::Int64 &value) const
+  {
+    return __sync_fetch_and_add(address,value);
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
+  {
+    return __sync_fetch_and_add(address,value);
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
+  {
+    return __sync_fetch_and_add(address,value);
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::UInt32 vtkmAtomicAdd(vtkm::Float32 *address, const vtkm::Float32 &value) const
+  {
+    return __sync_fetch_and_add(address,value);
+  }
+
+};
+
 }
 } // namespace vtkm::cont
 
diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h
index 3ca160e42..9a3499fe9 100644
--- a/vtkm/cont/testing/TestingDeviceAdapter.h
+++ b/vtkm/cont/testing/TestingDeviceAdapter.h
@@ -39,6 +39,8 @@
 
 #include <vtkm/cont/testing/Testing.h>
 
+#include <vtkm/exec/AtomicArray.h>
+
 #include <algorithm>
 #include <cmath>
 #include <utility>
@@ -276,6 +278,25 @@ public:
     }
   };
 
+  template<typename T>
+  struct AtomicKernel 
+  {
+    VTKM_CONT_EXPORT
+    AtomicKernel(const vtkm::exec::AtomicArray<T,DeviceAdapterTag> &array)
+    : AArray(array)
+    {  }
+
+    VTKM_EXEC_EXPORT void operator()(vtkm::Id index) const
+    {
+      this->AArray.Add(0, 1);
+    }
+
+    VTKM_CONT_EXPORT void SetErrorMessageBuffer(
+        const vtkm::exec::internal::ErrorMessageBuffer &) {  }
+
+    vtkm::exec::AtomicArray<T,DeviceAdapterTag> AArray;
+  };
+
 
 private:
 
@@ -1564,6 +1585,67 @@ private:
 
   }
 
+  static VTKM_CONT_EXPORT void TestAtomicArray()
+  {
+    // To test the atomics, ARRAY_SIZE number of threads will all increment  
+    // a single atomic value.
+    std::cout << "Testing Atomic Array with vtkm::Int32" << std::endl;
+    {
+    std::vector<vtkm::Int32> singleElement;
+    singleElement.push_back(0);
+    vtkm::cont::ArrayHandle<vtkm::Int32> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
+
+    vtkm::exec::AtomicArray<vtkm::Int32, DeviceAdapterTag> atomic(atomicElement);
+    Algorithm::Schedule(AtomicKernel<vtkm::Int32>(atomic), ARRAY_SIZE);
+    vtkm::Int32 expected = vtkm::Int32(ARRAY_SIZE);
+    vtkm::Int32 actual= atomicElement.GetPortalControl().Get(0);
+    std::cout<<"Atomic value "<<actual<<std::endl;
+    VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int32");
+    }
+
+    std::cout << "Testing Atomic Array with vtkm::UInt32" << std::endl;
+    {
+    std::vector<vtkm::UInt32> singleElement;
+    singleElement.push_back(0);
+    vtkm::cont::ArrayHandle<vtkm::UInt32> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
+
+    vtkm::exec::AtomicArray<vtkm::UInt32, DeviceAdapterTag> atomic(atomicElement);
+    Algorithm::Schedule(AtomicKernel<vtkm::UInt32>(atomic), ARRAY_SIZE);
+    vtkm::UInt32 expected = vtkm::UInt32(ARRAY_SIZE);
+    vtkm::UInt32 actual= atomicElement.GetPortalControl().Get(0);
+    std::cout<<"Atomic value "<<actual<<std::endl;
+    VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add UInt32");
+    }
+
+    std::cout << "Testing Atomic Array with vtkm::UInt64" << std::endl;
+    {
+    std::vector<vtkm::UInt64> singleElement;
+    singleElement.push_back(0);
+    vtkm::cont::ArrayHandle<vtkm::UInt64> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
+
+    vtkm::exec::AtomicArray<vtkm::UInt64, DeviceAdapterTag> atomic(atomicElement);
+    Algorithm::Schedule(AtomicKernel<vtkm::UInt64>(atomic), ARRAY_SIZE);
+    vtkm::UInt64 expected = vtkm::UInt64(ARRAY_SIZE);
+    vtkm::UInt64 actual= atomicElement.GetPortalControl().Get(0);
+    std::cout<<"Atomic value "<<actual<<std::endl;
+    VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add UInt64");
+    }
+
+    std::cout << "Testing Atomic Array with vtkm::Int64" << std::endl;
+    {
+    std::vector<vtkm::Int64> singleElement;
+    singleElement.push_back(0);
+    vtkm::cont::ArrayHandle<vtkm::Int64> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
+
+    vtkm::exec::AtomicArray<vtkm::Int64, DeviceAdapterTag> atomic(atomicElement);
+    Algorithm::Schedule(AtomicKernel<vtkm::Int64>(atomic), ARRAY_SIZE);
+    vtkm::Int64 expected = vtkm::Int64(ARRAY_SIZE);
+    vtkm::Int64 actual= atomicElement.GetPortalControl().Get(0);
+    std::cout<<"Atomic value "<<actual<<std::endl;
+    VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int64");
+    }
+  }
+
   struct TestAll
   {
     VTKM_CONT_EXPORT void operator()() const
@@ -1605,6 +1687,8 @@ private:
       TestStreamCompact();
 
       TestCopyArraysInDiffTypes();
+
+      TestAtomicArray();
     }
   };
 
diff --git a/vtkm/exec/AtomicArray.h b/vtkm/exec/AtomicArray.h
new file mode 100644
index 000000000..ff253e2eb
--- /dev/null
+++ b/vtkm/exec/AtomicArray.h
@@ -0,0 +1,63 @@
+//============================================================================
+//  Copyright (c) Kitware, Inc.
+//  All rights reserved.
+//  See LICENSE.txt for details.
+//  This software is distributed WITHOUT ANY WARRANTY; without even
+//  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+//  PURPOSE.  See the above copyright notice for more information.
+//
+//  Copyright 2015 Sandia Corporation.
+//  Copyright 2015 UT-Battelle, LLC.
+//  Copyright 2015 Los Alamos National Security.
+//
+//  Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
+//  the U.S. Government retains certain rights in this software.
+//
+//  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+//  Laboratory (LANL), the U.S. Government retains certain rights in
+//  this software.
+//============================================================================
+#ifndef vtk_m_exec_AtomicArray_h
+#define vtk_m_exec_AtomicArray_h
+
+#include <vtkm/cont/DeviceAdapter.h>
+#include <vtkm/cont/ArrayHandle.h>
+#include <vtkm/exec/ExecutionObjectBase.h>
+
+namespace vtkm {
+namespace exec {
+
+/// A class that can be used to atomically operate on an array of values
+/// safely across multiple instances of the same worklet. This is useful when
+/// you have an algorithm that needs to accumulate values in parallel, but writing
+/// out a value per worklet might be memory prohibitive.
+///
+/// To construct an AtomicArray you will need to pass in an vtkm::cont::ArrayHandle
+/// that is used as the underlying storage for the AtomicArray
+///
+template<typename T, typename DeviceAdapterTag>
+class AtomicArray : public vtkm::exec::ExecutionObjectBase
+{
+public:
+  template<typename StorageType>
+  VTKM_CONT_EXPORT
+  AtomicArray(vtkm::cont::ArrayHandle<T, StorageType> handle):
+    AtomicImplementation( handle )
+  {
+  }
+
+  VTKM_EXEC_EXPORT
+  T Add(vtkm::Id index, const T& value) const
+  {
+    return this->AtomicImplementation.Add(index,value);
+  }
+
+private:
+    vtkm::cont::DeviceAdapterAtomicArrayImplementation<T,DeviceAdapterTag>
+      AtomicImplementation;
+};
+
+}
+} // namespace vtkm::exec
+
+#endif //vtk_m_exec_AtomicArray_h
\ No newline at end of file

From ca37b2e9cbc13b10948bbc7475a60a76d04b154b Mon Sep 17 00:00:00 2001
From: mclarsen <larsen30@llnl.gov>
Date: Wed, 10 Feb 2016 08:21:38 -0800
Subject: [PATCH 02/27] Fixed compiler error with atomics

---
 vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h | 3 +++
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 6 ------
 vtkm/cont/testing/TestingDeviceAdapter.h           | 5 +----
 3 files changed, 4 insertions(+), 10 deletions(-)

diff --git a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
index 9ce72a3c6..ebc6a97fb 100644
--- a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
+++ b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
@@ -729,4 +729,7 @@ private:
   PortalType Portal;
 };
 
+}
+} // namespace vtkm::cont
+
 #endif //vtk_m_cont_internal_DeviceAdapterAlgorithmGeneral_h
diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 7d53d2cff..366ab6aea 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -349,12 +349,6 @@ private:
     return __sync_fetch_and_add(address,value);
   }
 
-  VTKM_EXEC_EXPORT
-  vtkm::UInt32 vtkmAtomicAdd(vtkm::Float32 *address, const vtkm::Float32 &value) const
-  {
-    return __sync_fetch_and_add(address,value);
-  }
-
 };
 
 }
diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h
index 9a3499fe9..5c492e7ae 100644
--- a/vtkm/cont/testing/TestingDeviceAdapter.h
+++ b/vtkm/cont/testing/TestingDeviceAdapter.h
@@ -1587,6 +1587,7 @@ private:
 
   static VTKM_CONT_EXPORT void TestAtomicArray()
   {
+    std::cout << "-------------------------------------------" << std::endl;
     // To test the atomics, ARRAY_SIZE number of threads will all increment  
     // a single atomic value.
     std::cout << "Testing Atomic Array with vtkm::Int32" << std::endl;
@@ -1599,7 +1600,6 @@ private:
     Algorithm::Schedule(AtomicKernel<vtkm::Int32>(atomic), ARRAY_SIZE);
     vtkm::Int32 expected = vtkm::Int32(ARRAY_SIZE);
     vtkm::Int32 actual= atomicElement.GetPortalControl().Get(0);
-    std::cout<<"Atomic value "<<actual<<std::endl;
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int32");
     }
 
@@ -1613,7 +1613,6 @@ private:
     Algorithm::Schedule(AtomicKernel<vtkm::UInt32>(atomic), ARRAY_SIZE);
     vtkm::UInt32 expected = vtkm::UInt32(ARRAY_SIZE);
     vtkm::UInt32 actual= atomicElement.GetPortalControl().Get(0);
-    std::cout<<"Atomic value "<<actual<<std::endl;
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add UInt32");
     }
 
@@ -1627,7 +1626,6 @@ private:
     Algorithm::Schedule(AtomicKernel<vtkm::UInt64>(atomic), ARRAY_SIZE);
     vtkm::UInt64 expected = vtkm::UInt64(ARRAY_SIZE);
     vtkm::UInt64 actual= atomicElement.GetPortalControl().Get(0);
-    std::cout<<"Atomic value "<<actual<<std::endl;
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add UInt64");
     }
 
@@ -1641,7 +1639,6 @@ private:
     Algorithm::Schedule(AtomicKernel<vtkm::Int64>(atomic), ARRAY_SIZE);
     vtkm::Int64 expected = vtkm::Int64(ARRAY_SIZE);
     vtkm::Int64 actual= atomicElement.GetPortalControl().Get(0);
-    std::cout<<"Atomic value "<<actual<<std::endl;
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int64");
     }
   }

From 78888e6542193220acab35638ddf3389dfa52b41 Mon Sep 17 00:00:00 2001
From: mclarsen <larsen30@llnl.gov>
Date: Wed, 2 Mar 2016 13:53:28 -0800
Subject: [PATCH 03/27] Testing Interlock exchange

---
 .../tbb/internal/DeviceAdapterAlgorithmTBB.h  | 19 +++++++++++++++++++
 1 file changed, 19 insertions(+)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 366ab6aea..c206f0ecb 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -67,6 +67,9 @@ VTKM_THIRDPARTY_PRE_INCLUDE
 #undef NOMINMAX
 #endif
 
+#if defined(VTKM_MSVC)
+#include <Windows.h>
+#endif
 VTKM_THIRDPARTY_POST_INCLUDE
 
 namespace vtkm {
@@ -328,25 +331,41 @@ private:
   VTKM_EXEC_EXPORT
   vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
   {
+#if defined(VTKM_MSVC)
+    return InterlockedExchangeAdd(address,value);
+#else
     return __sync_fetch_and_add(address,value);
+#endif
   }
 
   VTKM_EXEC_EXPORT
   vtkm::Int64 vtkmAtomicAdd(vtkm::Int64 *address, const vtkm::Int64 &value) const
   {
+#if defined(VTKM_MSVC)
+    return InterlockedExchangeAdd(address,value);
+#else
     return __sync_fetch_and_add(address,value);
+#endif
   }
 
   VTKM_EXEC_EXPORT
   vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
   {
+#if defined(VTKM_MSVC)
+    return InterlockedExchangeAdd64(address,value);
+#else
     return __sync_fetch_and_add(address,value);
+#endif
   }
 
   VTKM_EXEC_EXPORT
   vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
   {
+#if defined(VTKM_MSVC)
+    return InterlockedExchangeAdd64(address,value);
+#else
     return __sync_fetch_and_add(address,value);
+#endif
   }
 
 };

From 9a9f3a1d222bb12495bebe988f12565b5208b404 Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Wed, 2 Mar 2016 17:41:19 -0500
Subject: [PATCH 04/27] Forcing test


From f0c20903b3d041e2976af1ccfd10d2a72121a36f Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Thu, 3 Mar 2016 15:30:28 -0500
Subject: [PATCH 05/27] Sorting out iterator type

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index c206f0ecb..b65277427 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -316,8 +316,8 @@ public:
   VTKM_EXEC_EXPORT
   T Add(vtkm::Id index, const T& value) const
   {
-    T* lockedValue; 
-    lockedValue = (Iterators.GetBegin()+index);
+    IteratorType::pointer temp = &(*(Iterators.GetBegin()+index));
+    T* lockedValue = temp;
     return vtkmAtomicAdd(lockedValue, value);
   }
 
@@ -325,7 +325,7 @@ private:
   typedef typename vtkm::cont::ArrayHandle<T,vtkm::cont::StorageTagBasic>
         ::template ExecutionTypes<DeviceAdapterTagTBB>::Portal PortalType;
   typedef vtkm::cont::ArrayPortalToIterators<PortalType> IteratorsType;
-
+  typedef vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType IteratorType;
   IteratorsType Iterators;
 
   VTKM_EXEC_EXPORT

From 23e801da555691f3fcbbf6a0ffcf7b65cd93bda3 Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Thu, 3 Mar 2016 15:34:44 -0500
Subject: [PATCH 06/27] Adding typename

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index b65277427..83e680696 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -325,7 +325,7 @@ private:
   typedef typename vtkm::cont::ArrayHandle<T,vtkm::cont::StorageTagBasic>
         ::template ExecutionTypes<DeviceAdapterTagTBB>::Portal PortalType;
   typedef vtkm::cont::ArrayPortalToIterators<PortalType> IteratorsType;
-  typedef vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType IteratorType;
+  typedef typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType IteratorType;
   IteratorsType Iterators;
 
   VTKM_EXEC_EXPORT

From 1448d0d3b4d7fce6408c26002fa67bdbf9f99e2e Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Thu, 3 Mar 2016 16:04:04 -0500
Subject: [PATCH 07/27] Adding another typename

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 83e680696..9810eb727 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -316,7 +316,7 @@ public:
   VTKM_EXEC_EXPORT
   T Add(vtkm::Id index, const T& value) const
   {
-    IteratorType::pointer temp = &(*(Iterators.GetBegin()+index));
+    typename IteratorType::pointer temp = &(*(Iterators.GetBegin()+index));
     T* lockedValue = temp;
     return vtkmAtomicAdd(lockedValue, value);
   }

From 29ca4f38f6ec1fd3a539444230fd769500c30eba Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Thu, 3 Mar 2016 19:42:35 -0500
Subject: [PATCH 08/27] Attempting type match

---
 .../tbb/internal/DeviceAdapterAlgorithmTBB.h  | 25 ++++++++++++++-----
 1 file changed, 19 insertions(+), 6 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 9810eb727..137d8de51 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -316,8 +316,14 @@ public:
   VTKM_EXEC_EXPORT
   T Add(vtkm::Id index, const T& value) const
   {
+    T* lockedValue;
+#if defined(VTKM_MSVC)
+    typedef typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType IteratorType;
     typename IteratorType::pointer temp = &(*(Iterators.GetBegin()+index));
-    T* lockedValue = temp;
+    lockedValue = temp;
+#else
+    lockedValue = (Iterators.GetBegin()+index);
+#endif
     return vtkmAtomicAdd(lockedValue, value);
   }
 
@@ -325,14 +331,15 @@ private:
   typedef typename vtkm::cont::ArrayHandle<T,vtkm::cont::StorageTagBasic>
         ::template ExecutionTypes<DeviceAdapterTagTBB>::Portal PortalType;
   typedef vtkm::cont::ArrayPortalToIterators<PortalType> IteratorsType;
-  typedef typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType IteratorType;
   IteratorsType Iterators;
 
   VTKM_EXEC_EXPORT
   vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
   {
 #if defined(VTKM_MSVC)
-    return InterlockedExchangeAdd(address,value);
+    long msValue = value;
+    long * msPtr = (long *) address;
+    return InterlockedExchangeAdd(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);
 #endif
@@ -342,7 +349,9 @@ private:
   vtkm::Int64 vtkmAtomicAdd(vtkm::Int64 *address, const vtkm::Int64 &value) const
   {
 #if defined(VTKM_MSVC)
-    return InterlockedExchangeAdd(address,value);
+    long long msValue = value;
+    long long * msPtr = (long long *) address;
+    return InterlockedExchangeAdd64(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);
 #endif
@@ -352,7 +361,9 @@ private:
   vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
   {
 #if defined(VTKM_MSVC)
-    return InterlockedExchangeAdd64(address,value);
+    unsigned long msValue = value;
+    unsinged long * msPtr = (unsigned long *) address;
+    return InterlockedExchangeAdd(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);
 #endif
@@ -362,7 +373,9 @@ private:
   vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
   {
 #if defined(VTKM_MSVC)
-    return InterlockedExchangeAdd64(address,value);
+    unsigned long long msValue = value;
+    unsigned long long * msPtr = (unsigend long long *) address;
+    return InterlockedExchangeAdd64(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);
 #endif

From 4a28124658903613766797c66b11ae49a94ecfe7 Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Thu, 3 Mar 2016 22:17:29 -0500
Subject: [PATCH 09/27] Fixing typo

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 137d8de51..510a6adc9 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -362,7 +362,7 @@ private:
   {
 #if defined(VTKM_MSVC)
     unsigned long msValue = value;
-    unsinged long * msPtr = (unsigned long *) address;
+    unsigned long * msPtr = (unsigned long *) address;
     return InterlockedExchangeAdd(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);
@@ -374,7 +374,7 @@ private:
   {
 #if defined(VTKM_MSVC)
     unsigned long long msValue = value;
-    unsigned long long * msPtr = (unsigend long long *) address;
+    unsigned long long * msPtr = (unsignend long long *) address;
     return InterlockedExchangeAdd64(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);

From 25604a0288b8eca37eeb99b2d65ee6432e984e8d Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Thu, 3 Mar 2016 22:46:15 -0500
Subject: [PATCH 10/27] Another typo

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 510a6adc9..9f2af5dda 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -374,7 +374,7 @@ private:
   {
 #if defined(VTKM_MSVC)
     unsigned long long msValue = value;
-    unsigned long long * msPtr = (unsignend long long *) address;
+    unsigned long long * msPtr = (unsigned long long *) address;
     return InterlockedExchangeAdd64(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);

From 5df9088e1dfc76efa50b63ae0a5062de355a33f6 Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Thu, 3 Mar 2016 23:13:26 -0500
Subject: [PATCH 11/27] Adding volatile

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 9f2af5dda..59c3c0ea7 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -338,7 +338,7 @@ private:
   {
 #if defined(VTKM_MSVC)
     long msValue = value;
-    long * msPtr = (long *) address;
+    volatile long * msPtr = (volatile long *) address;
     return InterlockedExchangeAdd(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);
@@ -350,7 +350,7 @@ private:
   {
 #if defined(VTKM_MSVC)
     long long msValue = value;
-    long long * msPtr = (long long *) address;
+     volatile long long * msPtr = (volatile long long *) address;
     return InterlockedExchangeAdd64(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);
@@ -362,7 +362,7 @@ private:
   {
 #if defined(VTKM_MSVC)
     unsigned long msValue = value;
-    unsigned long * msPtr = (unsigned long *) address;
+    volatile unsigned long * msPtr = (volatile unsigned long *) address;
     return InterlockedExchangeAdd(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);
@@ -374,7 +374,7 @@ private:
   {
 #if defined(VTKM_MSVC)
     unsigned long long msValue = value;
-    unsigned long long * msPtr = (unsigned long long *) address;
+    volatile unsigned long long * msPtr = (volatile unsigned long long *) address;
     return InterlockedExchangeAdd64(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);

From f0f8f0b5566e96a4cf38f845777596ebf6263f79 Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Thu, 3 Mar 2016 23:36:21 -0500
Subject: [PATCH 12/27] attempt 6

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 59c3c0ea7..096b79253 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -373,8 +373,8 @@ private:
   vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
   {
 #if defined(VTKM_MSVC)
-    unsigned long long msValue = value;
-    volatile unsigned long long * msPtr = (volatile unsigned long long *) address;
+    long long msValue = value;
+    volatile long long * msPtr = (volatile long long *) address;
     return InterlockedExchangeAdd64(msPtr,msValue);
 #else
     return __sync_fetch_and_add(address,value);

From da66ea6a15f0e0990ee4e4b30517f2cf53acd45a Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Fri, 4 Mar 2016 10:18:35 -0800
Subject: [PATCH 13/27] Refactoring TBB atomics

---
 .../tbb/internal/DeviceAdapterAlgorithmTBB.h  | 63 ++++++++++---------
 1 file changed, 35 insertions(+), 28 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 096b79253..b76a25ca8 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -316,15 +316,17 @@ public:
   VTKM_EXEC_EXPORT
   T Add(vtkm::Id index, const T& value) const
   {
-    T* lockedValue;
 #if defined(VTKM_MSVC)
+    volatile T* lockedValue;
     typedef typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType IteratorType;
     typename IteratorType::pointer temp = &(*(Iterators.GetBegin()+index));
-    lockedValue = temp;
-#else
-    lockedValue = (Iterators.GetBegin()+index);
-#endif
+    lockedValue = static_cast<volatile T*>(temp);
     return vtkmAtomicAdd(lockedValue, value);
+#else
+    T* lockedValue;
+    lockedValue = (Iterators.GetBegin()+index);
+    return vtkmAtomicAdd(lockedValue, value);
+#endif
   }
 
 private:
@@ -333,50 +335,55 @@ private:
   typedef vtkm::cont::ArrayPortalToIterators<PortalType> IteratorsType;
   IteratorsType Iterators;
 
+#if defined(VTKM_MSVC) //MSVC atomics
   VTKM_EXEC_EXPORT
+  vtkm::Int32 vtkmAtomicAdd(volatile vtkm::Int32 *address, const vtkm::Int32 &value) const
+  {
+    return InterlockedExchangeAdd(address,value);
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::Int64 vtkmAtomicAdd(volatile vtkm::Int64 *address, const vtkm::Int64 &value) const
+  {
+    return InterlockedExchangeAdd64(address,value);
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::UInt32 vtkmAtomicAdd(volatile vtkm::UInt32 *address, const vtkm::UInt32 &value) const
+  {
+    return InterlockedExchangeAdd(address,value);
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::UInt64 vtkmAtomicAdd(volatile vtkm::UInt64 *address, const vtkm::UInt64 &value) const
+  {
+    return InterlockedExchangeAdd64(address,value);
+  }
+
+#else //gcc built-in atomics
+
+VTKM_EXEC_EXPORT
   vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
   {
-#if defined(VTKM_MSVC)
-    long msValue = value;
-    volatile long * msPtr = (volatile long *) address;
-    return InterlockedExchangeAdd(msPtr,msValue);
-#else
     return __sync_fetch_and_add(address,value);
-#endif
   }
 
   VTKM_EXEC_EXPORT
   vtkm::Int64 vtkmAtomicAdd(vtkm::Int64 *address, const vtkm::Int64 &value) const
   {
-#if defined(VTKM_MSVC)
-    long long msValue = value;
-     volatile long long * msPtr = (volatile long long *) address;
-    return InterlockedExchangeAdd64(msPtr,msValue);
-#else
     return __sync_fetch_and_add(address,value);
-#endif
   }
 
   VTKM_EXEC_EXPORT
   vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
   {
-#if defined(VTKM_MSVC)
-    unsigned long msValue = value;
-    volatile unsigned long * msPtr = (volatile unsigned long *) address;
-    return InterlockedExchangeAdd(msPtr,msValue);
-#else
     return __sync_fetch_and_add(address,value);
-#endif
   }
 
   VTKM_EXEC_EXPORT
   vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
   {
-#if defined(VTKM_MSVC)
-    long long msValue = value;
-    volatile long long * msPtr = (volatile long long *) address;
-    return InterlockedExchangeAdd64(msPtr,msValue);
-#else
+
     return __sync_fetch_and_add(address,value);
 #endif
   }

From 6d63961311b9e47f622c516b5ef924c43bb06706 Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Fri, 4 Mar 2016 10:42:07 -0800
Subject: [PATCH 14/27] Misplaced bracket

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index b76a25ca8..78f499b88 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -383,8 +383,8 @@ VTKM_EXEC_EXPORT
   VTKM_EXEC_EXPORT
   vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
   {
-
     return __sync_fetch_and_add(address,value);
+  }
 #endif
   }
 

From 47452d82bcdf819a6c3af8d57bb8db8706e8f121 Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Fri, 4 Mar 2016 10:49:43 -0800
Subject: [PATCH 15/27] fixing typo

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 78f499b88..9e0b74265 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -385,8 +385,8 @@ VTKM_EXEC_EXPORT
   {
     return __sync_fetch_and_add(address,value);
   }
+
 #endif
-  }
 
 };
 

From 733edc832524f988b3bbf1e52ace2f9e02843b92 Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Fri, 4 Mar 2016 11:13:53 -0800
Subject: [PATCH 16/27] changing MS method signatures to long instead of int

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 9e0b74265..4c60e1616 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -337,25 +337,25 @@ private:
 
 #if defined(VTKM_MSVC) //MSVC atomics
   VTKM_EXEC_EXPORT
-  vtkm::Int32 vtkmAtomicAdd(volatile vtkm::Int32 *address, const vtkm::Int32 &value) const
+  vtkm::Int32 vtkmAtomicAdd(volatile long *address, const long &value) const
   {
     return InterlockedExchangeAdd(address,value);
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::Int64 vtkmAtomicAdd(volatile vtkm::Int64 *address, const vtkm::Int64 &value) const
+  vtkm::Int64 vtkmAtomicAdd(volatile long long *address, const long long &value) const
   {
     return InterlockedExchangeAdd64(address,value);
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::UInt32 vtkmAtomicAdd(volatile vtkm::UInt32 *address, const vtkm::UInt32 &value) const
+  vtkm::UInt32 vtkmAtomicAdd(volatile unsigned long *address, const unsigned long &value) const
   {
     return InterlockedExchangeAdd(address,value);
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::UInt64 vtkmAtomicAdd(volatile vtkm::UInt64 *address, const vtkm::UInt64 &value) const
+  vtkm::UInt64 vtkmAtomicAdd(volatile unsigned long long *address, const unsigned long long &value) const
   {
     return InterlockedExchangeAdd64(address,value);
   }

From f3c7f188ee6a5b2ed5008a67b9f7a8b229f65cae Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Fri, 4 Mar 2016 11:52:52 -0800
Subject: [PATCH 17/27] Fixing problem with function signatures matching long *
 to int *

---
 .../tbb/internal/DeviceAdapterAlgorithmTBB.h  | 20 +++++++++----------
 1 file changed, 10 insertions(+), 10 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 4c60e1616..360ff8b9e 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -316,11 +316,11 @@ public:
   VTKM_EXEC_EXPORT
   T Add(vtkm::Id index, const T& value) const
   {
+    T* lockedValue;
 #if defined(VTKM_MSVC)
-    volatile T* lockedValue;
     typedef typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType IteratorType;
     typename IteratorType::pointer temp = &(*(Iterators.GetBegin()+index));
-    lockedValue = static_cast<volatile T*>(temp);
+    lockedValue = temp;
     return vtkmAtomicAdd(lockedValue, value);
 #else
     T* lockedValue;
@@ -337,27 +337,27 @@ private:
 
 #if defined(VTKM_MSVC) //MSVC atomics
   VTKM_EXEC_EXPORT
-  vtkm::Int32 vtkmAtomicAdd(volatile long *address, const long &value) const
+  vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
   {
-    return InterlockedExchangeAdd(address,value);
+    return InterlockedExchangeAdd(static_cast<volatile long *>(address),value);
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::Int64 vtkmAtomicAdd(volatile long long *address, const long long &value) const
+  vtkm::Int64 vtkmAtomicAdd(vtkm::Int64 *address, const vtkm::Int64 &value) const
   {
-    return InterlockedExchangeAdd64(address,value);
+    return InterlockedExchangeAdd64(static_cast<volatile long long *>(address),value);
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::UInt32 vtkmAtomicAdd(volatile unsigned long *address, const unsigned long &value) const
+  vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
   {
-    return InterlockedExchangeAdd(address,value);
+    return InterlockedExchangeAdd(static_cast<volatile unsigned long *>(address),value);
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::UInt64 vtkmAtomicAdd(volatile unsigned long long *address, const unsigned long long &value) const
+  vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
   {
-    return InterlockedExchangeAdd64(address,value);
+    return InterlockedExchangeAdd64(static_cast<volatile unsigned long long *>(address),value);
   }
 
 #else //gcc built-in atomics

From 72f2549bb90b6c6a6c5ea21b84fbf0ee968303e1 Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Fri, 4 Mar 2016 12:12:25 -0800
Subject: [PATCH 18/27] Removing shadow variable

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 1 -
 1 file changed, 1 deletion(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 360ff8b9e..163e323e3 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -323,7 +323,6 @@ public:
     lockedValue = temp;
     return vtkmAtomicAdd(lockedValue, value);
 #else
-    T* lockedValue;
     lockedValue = (Iterators.GetBegin()+index);
     return vtkmAtomicAdd(lockedValue, value);
 #endif

From ecc12dddb7887f908c2ae25a1a789a55c40a6c48 Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Fri, 4 Mar 2016 14:46:45 -0800
Subject: [PATCH 19/27] using reinterpret cast and removing warning from atomic
 array test

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h |  8 ++++----
 vtkm/cont/testing/TestingDeviceAdapter.h           | 13 ++++++++-----
 2 files changed, 12 insertions(+), 9 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 163e323e3..1ca663909 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -338,25 +338,25 @@ private:
   VTKM_EXEC_EXPORT
   vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
   {
-    return InterlockedExchangeAdd(static_cast<volatile long *>(address),value);
+    return InterlockedExchangeAdd(reinterpret_cast<volatile long *>(address),value);
   }
 
   VTKM_EXEC_EXPORT
   vtkm::Int64 vtkmAtomicAdd(vtkm::Int64 *address, const vtkm::Int64 &value) const
   {
-    return InterlockedExchangeAdd64(static_cast<volatile long long *>(address),value);
+    return InterlockedExchangeAdd64(reinterpret_cast<volatile long long *>(address),value);
   }
 
   VTKM_EXEC_EXPORT
   vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
   {
-    return InterlockedExchangeAdd(static_cast<volatile unsigned long *>(address),value);
+    return InterlockedExchangeAdd(reinterpret_cast<volatile unsigned long *>(address),value);
   }
 
   VTKM_EXEC_EXPORT
   vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
   {
-    return InterlockedExchangeAdd64(static_cast<volatile unsigned long long *>(address),value);
+    return InterlockedExchangeAdd64(reinterpret_cast<volatile unsigned long long *>(address),value);
   }
 
 #else //gcc built-in atomics
diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h
index 5c492e7ae..8e6d2906c 100644
--- a/vtkm/cont/testing/TestingDeviceAdapter.h
+++ b/vtkm/cont/testing/TestingDeviceAdapter.h
@@ -288,7 +288,8 @@ public:
 
     VTKM_EXEC_EXPORT void operator()(vtkm::Id index) const
     {
-      this->AArray.Add(0, 1);
+      T value = (T) index;
+      this->AArray.Add(0, value);
     }
 
     VTKM_CONT_EXPORT void SetErrorMessageBuffer(
@@ -1587,6 +1588,8 @@ private:
 
   static VTKM_CONT_EXPORT void TestAtomicArray()
   {
+    vtkm::Int32 atomicCount = 0;
+    for(vtkm::Int32 i = 0; i < ARRAY_SIZE; i++) atomicCount += i;
     std::cout << "-------------------------------------------" << std::endl;
     // To test the atomics, ARRAY_SIZE number of threads will all increment  
     // a single atomic value.
@@ -1598,7 +1601,7 @@ private:
 
     vtkm::exec::AtomicArray<vtkm::Int32, DeviceAdapterTag> atomic(atomicElement);
     Algorithm::Schedule(AtomicKernel<vtkm::Int32>(atomic), ARRAY_SIZE);
-    vtkm::Int32 expected = vtkm::Int32(ARRAY_SIZE);
+    vtkm::Int32 expected = vtkm::Int32(atomicCount);
     vtkm::Int32 actual= atomicElement.GetPortalControl().Get(0);
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int32");
     }
@@ -1611,7 +1614,7 @@ private:
 
     vtkm::exec::AtomicArray<vtkm::UInt32, DeviceAdapterTag> atomic(atomicElement);
     Algorithm::Schedule(AtomicKernel<vtkm::UInt32>(atomic), ARRAY_SIZE);
-    vtkm::UInt32 expected = vtkm::UInt32(ARRAY_SIZE);
+    vtkm::UInt32 expected = vtkm::UInt32(atomicCount);
     vtkm::UInt32 actual= atomicElement.GetPortalControl().Get(0);
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add UInt32");
     }
@@ -1624,7 +1627,7 @@ private:
 
     vtkm::exec::AtomicArray<vtkm::UInt64, DeviceAdapterTag> atomic(atomicElement);
     Algorithm::Schedule(AtomicKernel<vtkm::UInt64>(atomic), ARRAY_SIZE);
-    vtkm::UInt64 expected = vtkm::UInt64(ARRAY_SIZE);
+    vtkm::UInt64 expected = vtkm::UInt64(atomicCount);
     vtkm::UInt64 actual= atomicElement.GetPortalControl().Get(0);
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add UInt64");
     }
@@ -1637,7 +1640,7 @@ private:
 
     vtkm::exec::AtomicArray<vtkm::Int64, DeviceAdapterTag> atomic(atomicElement);
     Algorithm::Schedule(AtomicKernel<vtkm::Int64>(atomic), ARRAY_SIZE);
-    vtkm::Int64 expected = vtkm::Int64(ARRAY_SIZE);
+    vtkm::Int64 expected = vtkm::Int64(atomicCount);
     vtkm::Int64 actual= atomicElement.GetPortalControl().Get(0);
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int64");
     }

From 12632de34b592ea634d6204c4b9ff10cfbd476d9 Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Fri, 4 Mar 2016 15:17:56 -0800
Subject: [PATCH 20/27] Casting Uint64 to long long

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 1ca663909..a3ef8619a 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -356,7 +356,7 @@ private:
   VTKM_EXEC_EXPORT
   vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
   {
-    return InterlockedExchangeAdd64(reinterpret_cast<volatile unsigned long long *>(address),value);
+    return InterlockedExchangeAdd64(reinterpret_cast<volatile long long *>(address),value);
   }
 
 #else //gcc built-in atomics

From 3b46706e1f04e4f01a6c5c5a7ed22331b83e48ff Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Tue, 8 Mar 2016 09:41:02 -0800
Subject: [PATCH 21/27] Adding compare and swap and removing unsigned atomics

---
 .../internal/DeviceAdapterAlgorithmCuda.h     | 31 ++++---
 .../internal/DeviceAdapterAlgorithmGeneral.h  |  8 ++
 .../tbb/internal/DeviceAdapterAlgorithmTBB.h  | 33 +++++--
 vtkm/cont/testing/TestingDeviceAdapter.h      | 88 +++++++++++++------
 vtkm/exec/AtomicArray.h                       |  6 ++
 5 files changed, 115 insertions(+), 51 deletions(-)

diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
index 54cd62924..5ec7cd917 100644
--- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
+++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
@@ -205,6 +205,13 @@ public:
     return vtkmAtomicAdd(lockedValue, value);
   }
 
+  inline __device__
+  T CompareAndSwap(vtkm::Int64 *address, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
+  {
+    T *lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
+    return vtkmCompareAndSwap(lockedValue, value);
+  }
+
 private:
   typedef typename vtkm::cont::ArrayHandle<T,vtkm::cont::StorageTagBasic>
         ::template ExecutionTypes<vtkm::cont::DeviceAdapterTagCuda>::Portal PortalType;
@@ -216,18 +223,6 @@ private:
     return atomicAdd((unsigned long long *)address,(unsigned long long) value);
   }
 
-  inline __device__
-  vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
-  {
-    return atomicAdd((unsigned long long *)address,(unsigned long long) value);
-  }
-  
-  inline __device__
-  vtkm::Float32 vtkmAtomicAdd(vtkm::Float32 *address, const vtkm::Float32 &value) const
-  {
-    return atomicAdd(address,value);
-  }
-
   inline __device__
   vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
   {
@@ -235,9 +230,17 @@ private:
   }
 
   inline __device__
-  vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
+  vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32 *address, const vtkm::Int32 &newValue, const vtkm::Int32 &oldValue) const
   {
-    return atomicAdd(address,value);
+    return atomicCAS(address,oldValue,newValue);
+  }
+
+  inline __device__
+  vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
+  {
+    return atomicCAS(static_cast<vtkm::UInt64*>(address),
+                     static_cast<vtkm::UInt64>(oldValue),
+                     static_cast<vtkm::UInt64>(newValue));
   }  
 };
 
diff --git a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
index ebc6a97fb..06e8b4952 100644
--- a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
+++ b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
@@ -723,6 +723,14 @@ public:
     return old;
   }
 
+  VTKM_EXEC_EXPORT
+  T CompareAndSwap(vtkm::Id index, const T& newValue, const T& oldValue) const
+  {
+    const T old = this->Portal.Get(index);
+    if(old == oldValue) this->Portal.Set(index,newValue);
+    return old;
+  }
+
 private:
     typedef typename vtkm::cont::ArrayHandle<T>
               ::template ExecutionTypes<DeviceTag>::Portal PortalType;
diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index a3ef8619a..1dd8bd53e 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -328,6 +328,21 @@ public:
 #endif
   }
 
+  VTKM_EXEC_EXPORT
+  T CompareAndSwap(vtkm::Id index, const T& newValue, const T& oldValue) const
+  {
+    T* lockedValue;
+#if defined(VTKM_MSVC)
+    typedef typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType IteratorType;
+    typename IteratorType::pointer temp = &(*(Iterators.GetBegin()+index));
+    lockedValue = temp;
+    return vtkmCompareAndSwap(lockedValue, newValue, oldValue);
+#else
+    lockedValue = (Iterators.GetBegin()+index);
+    return vtkmCompareAndSwap(lockedValue, newValue, value);
+#endif
+  }
+
 private:
   typedef typename vtkm::cont::ArrayHandle<T,vtkm::cont::StorageTagBasic>
         ::template ExecutionTypes<DeviceAdapterTagTBB>::Portal PortalType;
@@ -348,20 +363,20 @@ private:
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
+  vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32 *address, const vtkm::Int32 &newValue, const vtkm::Int32 &oldValue) const
   {
-    return InterlockedExchangeAdd(reinterpret_cast<volatile unsigned long *>(address),value);
+    return InterlockedExchangeAdd(reinterpret_cast<volatile long *>(address),newValue,oldValue);
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
+  vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address,const vtkm::Int64 &newValue const vtkm::Int64 &oldValue) const
   {
-    return InterlockedExchangeAdd64(reinterpret_cast<volatile long long *>(address),value);
+    return InterlockedExchangeAdd64(reinterpret_cast<volatile long long *>(address),newValue, oldValue);
   }
 
 #else //gcc built-in atomics
 
-VTKM_EXEC_EXPORT
+  VTKM_EXEC_EXPORT
   vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
   {
     return __sync_fetch_and_add(address,value);
@@ -374,15 +389,15 @@ VTKM_EXEC_EXPORT
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
+  vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32 *address, const vtkm::Int32 &newValue, const vtkm::Int32 &oldValue) const
   {
-    return __sync_fetch_and_add(address,value);
+    return __sync_val_compare_and_swap(address,oldValue, newValue);
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::UInt64 vtkmAtomicAdd(vtkm::UInt64 *address, const vtkm::UInt64 &value) const
+  vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address,const vtkm::Int64 &newValue const vtkm::Int64 &oldValue) const
   {
-    return __sync_fetch_and_add(address,value);
+    return __sync_val_compare_and_swap(address,oldValue,newValue);
   }
 
 #endif
diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h
index 8e6d2906c..60dbcaf90 100644
--- a/vtkm/cont/testing/TestingDeviceAdapter.h
+++ b/vtkm/cont/testing/TestingDeviceAdapter.h
@@ -298,6 +298,36 @@ public:
     vtkm::exec::AtomicArray<T,DeviceAdapterTag> AArray;
   };
 
+  template<typename T>
+  struct AtomicCASKernel 
+  {
+    VTKM_CONT_EXPORT
+    AtomicCASKernel(const vtkm::exec::AtomicArray<T,DeviceAdapterTag> &array)
+    : AArray(array)
+    {  }
+
+    VTKM_EXEC_EXPORT void operator()(vtkm::Id index) const
+    {
+      T value = (T) index;
+      //Get the old value from the array with a no-op
+      T oldValue = this->AArray.Add(0,T(0));
+      //This creates an atomic add using the CAS operatoin
+      T assumed = T(0);
+      do
+      { 
+        assumed = oldValue;
+        oldValue = this->AArray.CompareAndSwap(0, (assumed + value) , assumed);
+
+      } while (assumed != oldValue);
+      
+    }
+
+    VTKM_CONT_EXPORT void SetErrorMessageBuffer(
+        const vtkm::exec::internal::ErrorMessageBuffer &) {  }
+
+    vtkm::exec::AtomicArray<T,DeviceAdapterTag> AArray;
+  };
+
 
 private:
 
@@ -1593,7 +1623,7 @@ private:
     std::cout << "-------------------------------------------" << std::endl;
     // To test the atomics, ARRAY_SIZE number of threads will all increment  
     // a single atomic value.
-    std::cout << "Testing Atomic Array with vtkm::Int32" << std::endl;
+    std::cout << "Testing Atomic Add with vtkm::Int32" << std::endl;
     {
     std::vector<vtkm::Int32> singleElement;
     singleElement.push_back(0);
@@ -1606,33 +1636,7 @@ private:
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int32");
     }
 
-    std::cout << "Testing Atomic Array with vtkm::UInt32" << std::endl;
-    {
-    std::vector<vtkm::UInt32> singleElement;
-    singleElement.push_back(0);
-    vtkm::cont::ArrayHandle<vtkm::UInt32> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
-
-    vtkm::exec::AtomicArray<vtkm::UInt32, DeviceAdapterTag> atomic(atomicElement);
-    Algorithm::Schedule(AtomicKernel<vtkm::UInt32>(atomic), ARRAY_SIZE);
-    vtkm::UInt32 expected = vtkm::UInt32(atomicCount);
-    vtkm::UInt32 actual= atomicElement.GetPortalControl().Get(0);
-    VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add UInt32");
-    }
-
-    std::cout << "Testing Atomic Array with vtkm::UInt64" << std::endl;
-    {
-    std::vector<vtkm::UInt64> singleElement;
-    singleElement.push_back(0);
-    vtkm::cont::ArrayHandle<vtkm::UInt64> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
-
-    vtkm::exec::AtomicArray<vtkm::UInt64, DeviceAdapterTag> atomic(atomicElement);
-    Algorithm::Schedule(AtomicKernel<vtkm::UInt64>(atomic), ARRAY_SIZE);
-    vtkm::UInt64 expected = vtkm::UInt64(atomicCount);
-    vtkm::UInt64 actual= atomicElement.GetPortalControl().Get(0);
-    VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add UInt64");
-    }
-
-    std::cout << "Testing Atomic Array with vtkm::Int64" << std::endl;
+    std::cout << "Testing Atomic Add with vtkm::Int64" << std::endl;
     {
     std::vector<vtkm::Int64> singleElement;
     singleElement.push_back(0);
@@ -1644,6 +1648,34 @@ private:
     vtkm::Int64 actual= atomicElement.GetPortalControl().Get(0);
     VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic add Int64");
     }
+
+    std::cout << "Testing Atomic CAS with vtkm::Int32" << std::endl;
+    {
+    std::vector<vtkm::Int32> singleElement;
+    singleElement.push_back(0);
+    vtkm::cont::ArrayHandle<vtkm::Int32> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
+
+    vtkm::exec::AtomicArray<vtkm::Int32, DeviceAdapterTag> atomic(atomicElement);
+    Algorithm::Schedule(AtomicCASKernel<vtkm::Int32>(atomic), ARRAY_SIZE);
+    vtkm::Int32 expected = vtkm::Int32(atomicCount);
+    vtkm::Int32 actual= atomicElement.GetPortalControl().Get(0);
+    VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic CAS Int32");
+    }
+
+    std::cout << "Testing Atomic CAS with vtkm::Int64" << std::endl;
+    {
+    std::vector<vtkm::Int64> singleElement;
+    singleElement.push_back(0);
+    vtkm::cont::ArrayHandle<vtkm::Int64> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
+
+    vtkm::exec::AtomicArray<vtkm::Int64, DeviceAdapterTag> atomic(atomicElement);
+    Algorithm::Schedule(AtomicCASKernel<vtkm::Int64>(atomic), ARRAY_SIZE);
+    vtkm::Int64 expected = vtkm::Int64(atomicCount);
+    vtkm::Int64 actual= atomicElement.GetPortalControl().Get(0);
+    VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic CAS Int64");
+    }
+
+    
   }
 
   struct TestAll
diff --git a/vtkm/exec/AtomicArray.h b/vtkm/exec/AtomicArray.h
index ff253e2eb..523f3a5ae 100644
--- a/vtkm/exec/AtomicArray.h
+++ b/vtkm/exec/AtomicArray.h
@@ -52,6 +52,12 @@ public:
     return this->AtomicImplementation.Add(index,value);
   }
 
+  VTKM_EXEC_EXPORT
+  T CompareAndSwap(vtkm::Id index, const T& newValue, const T& oldValue) const
+  {
+    return this->AtomicImplementation.CompareAndSwap(index,newValue, oldValue);
+  }
+
 private:
     vtkm::cont::DeviceAdapterAtomicArrayImplementation<T,DeviceAdapterTag>
       AtomicImplementation;

From 40b6db7eee0bb17d791de921922a117a73740498 Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Tue, 8 Mar 2016 09:51:50 -0800
Subject: [PATCH 22/27] Inserted missing ,

---
 vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h | 2 +-
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h   | 4 ++--
 2 files changed, 3 insertions(+), 3 deletions(-)

diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
index 5ec7cd917..e8af4b40b 100644
--- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
+++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
@@ -209,7 +209,7 @@ public:
   T CompareAndSwap(vtkm::Int64 *address, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
   {
     T *lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
-    return vtkmCompareAndSwap(lockedValue, value);
+    return vtkmCompareAndSwap(lockedValue, oldValue);
   }
 
 private:
diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 1dd8bd53e..5e6c27634 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -339,7 +339,7 @@ public:
     return vtkmCompareAndSwap(lockedValue, newValue, oldValue);
 #else
     lockedValue = (Iterators.GetBegin()+index);
-    return vtkmCompareAndSwap(lockedValue, newValue, value);
+    return vtkmCompareAndSwap(lockedValue, newValue, oldValue);
 #endif
   }
 
@@ -395,7 +395,7 @@ private:
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address,const vtkm::Int64 &newValue const vtkm::Int64 &oldValue) const
+  vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address,const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
   {
     return __sync_val_compare_and_swap(address,oldValue,newValue);
   }

From 43131ee02b5968afaf0caed71093724b4c791316 Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Tue, 8 Mar 2016 09:58:20 -0800
Subject: [PATCH 23/27] Adding comments about CAS

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 2 +-
 vtkm/exec/AtomicArray.h                            | 5 +++++
 2 files changed, 6 insertions(+), 1 deletion(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 5e6c27634..988ba421b 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -369,7 +369,7 @@ private:
   }
 
   VTKM_EXEC_EXPORT
-  vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address,const vtkm::Int64 &newValue const vtkm::Int64 &oldValue) const
+  vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address,const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
   {
     return InterlockedExchangeAdd64(reinterpret_cast<volatile long long *>(address),newValue, oldValue);
   }
diff --git a/vtkm/exec/AtomicArray.h b/vtkm/exec/AtomicArray.h
index 523f3a5ae..072ef84ae 100644
--- a/vtkm/exec/AtomicArray.h
+++ b/vtkm/exec/AtomicArray.h
@@ -52,6 +52,11 @@ public:
     return this->AtomicImplementation.Add(index,value);
   }
 
+  //
+  // Compare and Swap is an atomic exchange operation. If the value at 
+  // the index is equal to oldValue, then newValue is written to the index.
+  // The operation was successful if return value is eqaul to oldValue
+  //
   VTKM_EXEC_EXPORT
   T CompareAndSwap(vtkm::Id index, const T& newValue, const T& oldValue) const
   {

From 249cce352bc2d55f5c870045ab4954e3b6ef98ca Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Tue, 8 Mar 2016 10:39:23 -0800
Subject: [PATCH 24/27] Adding type restrictions to serial atomics

---
 .../internal/DeviceAdapterAlgorithmCuda.h     |  8 ++--
 .../internal/DeviceAdapterAlgorithmGeneral.h  | 41 ++++++++++++++++---
 2 files changed, 39 insertions(+), 10 deletions(-)

diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
index e8af4b40b..17adde2f9 100644
--- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
+++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
@@ -209,7 +209,7 @@ public:
   T CompareAndSwap(vtkm::Int64 *address, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
   {
     T *lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
-    return vtkmCompareAndSwap(lockedValue, oldValue);
+    return vtkmCompareAndSwap(lockedValue, newValue, oldValue);
   }
 
 private:
@@ -238,9 +238,9 @@ private:
   inline __device__
   vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
   {
-    return atomicCAS(static_cast<vtkm::UInt64*>(address),
-                     static_cast<vtkm::UInt64>(oldValue),
-                     static_cast<vtkm::UInt64>(newValue));
+    return atomicCAS((unsigned long long int*) address,
+                     (unsigned long long int ) oldValue,
+                     (unsigned long long int ) newValue);
   }  
 };
 
diff --git a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
index 06e8b4952..40db13fe5 100644
--- a/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
+++ b/vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
@@ -718,23 +718,52 @@ public:
   VTKM_EXEC_EXPORT
   T Add(vtkm::Id index, const T& value) const
   {
-    const T old = this->Portal.Get(index);
-    this->Portal.Set(index, old + value);
-    return old;
+    return vtkmAtomicAdd(index, value);
   }
 
   VTKM_EXEC_EXPORT
   T CompareAndSwap(vtkm::Id index, const T& newValue, const T& oldValue) const
   {
-    const T old = this->Portal.Get(index);
-    if(old == oldValue) this->Portal.Set(index,newValue);
-    return old;
+    return vtkmCompareAndSwap(index, newValue, oldValue);
   }
 
 private:
     typedef typename vtkm::cont::ArrayHandle<T>
               ::template ExecutionTypes<DeviceTag>::Portal PortalType;
   PortalType Portal;
+
+  VTKM_EXEC_EXPORT
+  vtkm::Int32 vtkmAtomicAdd(const vtkm::Id &index, const vtkm::Int32 &value) const
+  {
+    const vtkm::Int32 old = this->Portal.Get(index);
+    this->Portal.Set(index, old + value);
+    return old;
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::Int64 vtkmAtomicAdd(const vtkm::Id &index, const vtkm::Int64 &value) const
+  {
+    const vtkm::Int64 old = this->Portal.Get(index);
+    this->Portal.Set(index, old + value);
+    return old;
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::Int32 vtkmCompareAndSwap(const vtkm::Id &index, const vtkm::Int32 &newValue, const vtkm::Int32 &oldValue) const
+  {
+    const vtkm::Int32 old = this->Portal.Get(index);
+    if(old == oldValue) this->Portal.Set(index, newValue);
+    return old;
+  }
+
+  VTKM_EXEC_EXPORT
+  vtkm::Int64 vtkmCompareAndSwap(const vtkm::Id &index, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
+  {
+    const vtkm::Int64 old = this->Portal.Get(index);
+    if(old == oldValue) this->Portal.Set(index, newValue);
+    return old;
+  }
+
 };
 
 }

From 00a7f6c14b9224b7f997de2aa9f4cf4b1e495f3c Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Tue, 8 Mar 2016 12:04:55 -0800
Subject: [PATCH 25/27] Correcting function call

---
 vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
index 988ba421b..ead428c09 100644
--- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
+++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h
@@ -365,13 +365,13 @@ private:
   VTKM_EXEC_EXPORT
   vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32 *address, const vtkm::Int32 &newValue, const vtkm::Int32 &oldValue) const
   {
-    return InterlockedExchangeAdd(reinterpret_cast<volatile long *>(address),newValue,oldValue);
+    return InterlockedCompareExchange(reinterpret_cast<volatile long *>(address),newValue,oldValue);
   }
 
   VTKM_EXEC_EXPORT
   vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address,const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
   {
-    return InterlockedExchangeAdd64(reinterpret_cast<volatile long long *>(address),newValue, oldValue);
+    return InterlockedCompareExchange64(reinterpret_cast<volatile long long *>(address),newValue, oldValue);
   }
 
 #else //gcc built-in atomics

From e5c4aa3f78d140efc787d37e044a686dee939a6b Mon Sep 17 00:00:00 2001
From: Matt Larsen <larsen30@llnl.gov>
Date: Tue, 8 Mar 2016 12:41:11 -0800
Subject: [PATCH 26/27] Fixing cuda index error

---
 vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
index 17adde2f9..22a65e65c 100644
--- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
+++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
@@ -206,7 +206,7 @@ public:
   }
 
   inline __device__
-  T CompareAndSwap(vtkm::Int64 *address, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
+  T CompareAndSwap(vtkm::Id index, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
   {
     T *lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
     return vtkmCompareAndSwap(lockedValue, newValue, oldValue);

From 5ddade7a44f6635acf5da892aeec3162cd8a5803 Mon Sep 17 00:00:00 2001
From: Matt Larsen <mlarsen@cs.uoregon.edu>
Date: Wed, 9 Mar 2016 14:29:59 -0500
Subject: [PATCH 27/27] Adding some basic documentation on atomics.

---
 vtkm/exec/AtomicArray.h | 5 +++++
 1 file changed, 5 insertions(+)

diff --git a/vtkm/exec/AtomicArray.h b/vtkm/exec/AtomicArray.h
index 072ef84ae..51b9de794 100644
--- a/vtkm/exec/AtomicArray.h
+++ b/vtkm/exec/AtomicArray.h
@@ -35,6 +35,11 @@ namespace exec {
 /// To construct an AtomicArray you will need to pass in an vtkm::cont::ArrayHandle
 /// that is used as the underlying storage for the AtomicArray
 ///
+/// Supported Operations: add / compare and swap (CAS)
+///                     
+/// Supported Types: 32 / 64 bit signed integers
+///                 
+///
 template<typename T, typename DeviceAdapterTag>
 class AtomicArray : public vtkm::exec::ExecutionObjectBase
 {