Updates will be applied April 15th at 12pm EDT (UTC-0400). GitLab could be a little slow between 12 - 12:45pm EDT.

Commit 3b46706e authored by Matt Larsen's avatar Matt Larsen

Adding compare and swap and removing unsigned atomics

parent 12632de3
......@@ -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;
......@@ -217,27 +224,23 @@ private:
}
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
vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &value) const
{
return atomicAdd(address,value);
}
inline __device__
vtkm::Int32 vtkmAtomicAdd(vtkm::Int32 *address, const vtkm::Int32 &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::UInt32 vtkmAtomicAdd(vtkm::UInt32 *address, const vtkm::UInt32 &value) const
vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64 *address, const vtkm::Int64 &newValue, const vtkm::Int64 &oldValue) const
{
return atomicAdd(address,value);
return atomicCAS(static_cast<vtkm::UInt64*>(address),
static_cast<vtkm::UInt64>(oldValue),
static_cast<vtkm::UInt64>(newValue));
}
};
......
......@@ -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;
......
......@@ -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
......
......@@ -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,44 +1636,46 @@ 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::cout << "Testing Atomic Add with vtkm::Int64" << std::endl;
{
std::vector<vtkm::UInt32> singleElement;
std::vector<vtkm::Int64> singleElement;
singleElement.push_back(0);
vtkm::cont::ArrayHandle<vtkm::UInt32> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
vtkm::cont::ArrayHandle<vtkm::Int64> 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");
vtkm::exec::AtomicArray<vtkm::Int64, DeviceAdapterTag> atomic(atomicElement);
Algorithm::Schedule(AtomicKernel<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 add Int64");
}
std::cout << "Testing Atomic Array with vtkm::UInt64" << std::endl;
std::cout << "Testing Atomic CAS with vtkm::Int32" << std::endl;
{
std::vector<vtkm::UInt64> singleElement;
std::vector<vtkm::Int32> singleElement;
singleElement.push_back(0);
vtkm::cont::ArrayHandle<vtkm::UInt64> atomicElement = vtkm::cont::make_ArrayHandle(singleElement);
vtkm::cont::ArrayHandle<vtkm::Int32> 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");
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 Array with vtkm::Int64" << std::endl;
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(AtomicKernel<vtkm::Int64>(atomic), ARRAY_SIZE);
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 add Int64");
VTKM_TEST_ASSERT(expected == actual, "Did not get expected value: Atomic CAS Int64");
}
}
struct TestAll
......
......@@ -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;
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment