Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-16 17:22:55 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Adding type restrictions to serial atomics

Browse Source
This commit is contained in:
Matt Larsen 2016-03-08 10:39:23 -08:00
parent 43131ee02b
commit 249cce352b
2 changed files with 39 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
View File

@ -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);
}
};

41
vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h
View File

@ -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;
}
};
}