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

Update atomic interfaces to support Add/CAS for UInt32/64.

Browse Source 
These will be used for the AtomicArray implementation.
This commit is contained in:
Allison Vacanti 2019-08-22 15:45:39 -04:00
parent 720b452eb4
commit 0e728c8000
3 changed files with 76 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

40
vtkm/cont/cuda/internal/AtomicInterfaceExecutionCuda.h
View File

@ -50,6 +50,13 @@ public:
__threadfence(); \
*vaddr = value; \
} \
VTKM_SUPPRESS_EXEC_WARNINGS __device__ static type Add(type* addr, type arg) \
{ \
__threadfence(); \
auto result = atomicAdd(addr, arg); \
__threadfence(); \
return result; \
} \
VTKM_SUPPRESS_EXEC_WARNINGS __device__ static type Not(type* addr) \
{ \
return AtomicInterfaceExecution::Xor(addr, static_cast<type>(~type{ 0u })); \
@ -87,6 +94,39 @@ public:
VTKM_ATOMIC_OPS_FOR_TYPE(vtkm::UInt32)
#undef VTKM_ATOMIC_OPS_FOR_TYPE
// We also support Load, Add & CAS for 64-bit unsigned ints in order to
// support AtomicArray usecases. We can't generally support UInt64 without
// bumping our minimum device req to compute capability 3.5 (though we could
// just use CAS for everything if this becomes a need). All of our supported
// devices do support add / CAS on UInt64, just not all the bit stuff.
VTKM_SUPPRESS_EXEC_WARNINGS __device__ static vtkm::UInt64 Load(const vtkm::UInt64* addr)
{
const volatile vtkm::UInt64* vaddr = addr; /* volatile to bypass cache*/
const vtkm::UInt64 value = *vaddr;
/* fence to ensure that dependent reads are correctly ordered */
__threadfence();
return value;
}
VTKM_SUPPRESS_EXEC_WARNINGS __device__ static vtkm::UInt64 Add(vtkm::UInt64* addr,
vtkm::UInt64 arg)
{
__threadfence();
auto result = atomicAdd(addr, arg);
__threadfence();
return result;
}
VTKM_SUPPRESS_EXEC_WARNINGS __device__ static vtkm::UInt64 CompareAndSwap(vtkm::UInt64* addr,
vtkm::UInt64 newWord,
vtkm::UInt64 expected)
{
__threadfence();
auto result = atomicCAS(addr, expected, newWord);
__threadfence();
return result;
}
};
}
}

15
vtkm/cont/internal/AtomicInterfaceControl.h
View File

@ -56,6 +56,12 @@ private:
return dst;
}
template <typename T>
VTKM_EXEC_CONT static T BitCast(T&& src)
{
return std::forward<T>(src);
}
public:
// Note about Load and Store implementations:
//
@ -127,6 +133,11 @@ public:
}
#define VTKM_ATOMIC_OPS_FOR_TYPE(vtkmType, winType, suffix) \
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_EXEC_CONT static vtkmType Add(vtkmType* addr, vtkmType arg) \
{ \
return BitCast<vtkmType>(_InterlockedExchangeAdd##suffix( \
reinterpret_cast<volatile winType*>(addr), BitCast<winType>(arg))); \
} \
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_EXEC_CONT static vtkmType Not(vtkmType* addr) \
{ \
return Xor(addr, static_cast<vtkmType>(~vtkmType{ 0u })); \
@ -173,6 +184,10 @@ public:
{ \
return __atomic_store_n(addr, value, __ATOMIC_RELEASE); \
} \
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_EXEC_CONT static type Add(type* addr, type arg) \
{ \
return __atomic_fetch_add(addr, arg, __ATOMIC_SEQ_CST); \
} \
VTKM_SUPPRESS_EXEC_WARNINGS VTKM_EXEC_CONT static type Not(type* addr) \
{ \
return Xor(addr, static_cast<type>(~type{ 0u })); \

21
vtkm/cont/internal/AtomicInterfaceExecution.h
View File

@ -29,6 +29,9 @@ namespace internal
/// be the same. A full list of supported word types is advertised in the type
/// list @a WordTypes.
///
/// In addition, each atomic interface must support Add and CompareAndSwap on
/// UInt32 and UInt64, as these are required for the AtomicArray implementation.
///
/// To implement this on devices that share the control environment, subclass
/// vtkm::cont::internal::AtomicInterfaceControl, which may also be used
/// directly from control-side code.
@ -52,6 +55,19 @@ class AtomicInterfaceExecution
VTKM_EXEC static void Store(vtkm::WordTypeDefault* addr, vtkm::WordTypeDefault value);
VTKM_EXEC static void Store(WordTypePreferred* addr, WordTypePreferred value);
/// Perform an atomic integer add operation on the word at @a addr, adding
/// @arg. This operation performs a full memory barrier around the atomic
/// access.
///
/// The value at @a addr prior to the addition is returned.
///
/// @note Overflow behavior is not defined for this operation.
/// @{
VTKM_EXEC static vtkm::WordTypeDefault Add(vtkm::WordTypeDefault* addr,
vtkm::WordTypeDefault arg);
VTKM_EXEC static WordTypePreferred Add(WordTypePreferred* addr, WordTypePreferred arg);
/// @}
/// Perform a bitwise atomic not operation on the word at @a addr.
/// This operation performs a full memory barrier around the atomic access.
/// @{
@ -84,6 +100,11 @@ class AtomicInterfaceExecution
/// @}
/// Perform an atomic CAS operation on the word at @a addr.
///
/// If the value at @a addr equals @a expected, @a addr will be set to
/// @a newWord and @a expected is returned. Otherwise, the value at @a addr
/// is returned and not modified.
///
/// This operation performs a full memory barrier around the atomic access.
/// @{
VTKM_EXEC static vtkm::WordTypeDefault CompareAndSwap(vtkm::WordTypeDefault* addr,