Move DevAdaptAtomicArrayImplementation to its own file.
This commit is contained in:
parent
3af9f66083
commit
be0c6a17a9
@ -23,6 +23,7 @@ set(headers
|
||||
CudaAllocator.h
|
||||
DeviceAdapterAlgorithmCuda.h
|
||||
DeviceAdapterAlgorithmThrust.h
|
||||
DeviceAdapterAtomicArrayImplementationCuda.h
|
||||
DeviceAdapterRuntimeDetectorCuda.h
|
||||
DeviceAdapterTagCuda.h
|
||||
DeviceAdapterTimerImplementationCuda.h
|
||||
|
@ -56,66 +56,6 @@ struct DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagCuda>
|
||||
}
|
||||
};
|
||||
|
||||
/// CUDA contains its own atomic operations
|
||||
///
|
||||
template <typename T>
|
||||
class DeviceAdapterAtomicArrayImplementation<T, vtkm::cont::DeviceAdapterTagCuda>
|
||||
{
|
||||
public:
|
||||
VTKM_CONT
|
||||
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);
|
||||
}
|
||||
|
||||
inline __device__ 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);
|
||||
}
|
||||
|
||||
private:
|
||||
using PortalType =
|
||||
typename vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>::template ExecutionTypes<
|
||||
vtkm::cont::DeviceAdapterTagCuda>::Portal;
|
||||
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::Int32 vtkmAtomicAdd(vtkm::Int32* address, const vtkm::Int32& value) const
|
||||
{
|
||||
return atomicAdd(address, value);
|
||||
}
|
||||
|
||||
inline __device__ vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
|
||||
const vtkm::Int32& newValue,
|
||||
const vtkm::Int32& oldValue) const
|
||||
{
|
||||
return atomicCAS(address, oldValue, newValue);
|
||||
}
|
||||
|
||||
inline __device__ vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
|
||||
const vtkm::Int64& newValue,
|
||||
const vtkm::Int64& oldValue) const
|
||||
{
|
||||
return atomicCAS((unsigned long long int*)address,
|
||||
(unsigned long long int)oldValue,
|
||||
(unsigned long long int)newValue);
|
||||
}
|
||||
};
|
||||
|
||||
template <>
|
||||
class DeviceTaskTypes<vtkm::cont::DeviceAdapterTagCuda>
|
||||
{
|
||||
|
@ -32,6 +32,7 @@
|
||||
|
||||
#include <vtkm/cont/cuda/ErrorCuda.h>
|
||||
|
||||
#include <vtkm/cont/cuda/internal/DeviceAdapterAtomicArrayImplementationCuda.h>
|
||||
#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
|
||||
#include <vtkm/cont/cuda/internal/MakeThrustIterator.h>
|
||||
#include <vtkm/cont/cuda/internal/ThrustExceptionHandler.h>
|
||||
|
@ -0,0 +1,104 @@
|
||||
//============================================================================
|
||||
// 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 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
|
||||
// Copyright 2018 UT-Battelle, LLC.
|
||||
// Copyright 2018 Los Alamos National Security.
|
||||
//
|
||||
// Under the terms of Contract DE-NA0003525 with NTESS,
|
||||
// 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_cont_internal_DeviceAdapterAtomicArrayImplementationCuda_h
|
||||
#define vtk_m_cont_internal_DeviceAdapterAtomicArrayImplementationCuda_h
|
||||
|
||||
#include <vtkm/Types.h>
|
||||
|
||||
#include <vtkm/cont/ArrayHandle.h>
|
||||
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
|
||||
#include <vtkm/cont/StorageBasic.h>
|
||||
|
||||
#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
|
||||
|
||||
// Disable warnings we check vtkm for but Thrust does not.
|
||||
VTKM_THIRDPARTY_PRE_INCLUDE
|
||||
#include <thrust/device_ptr.h>
|
||||
VTKM_THIRDPARTY_POST_INCLUDE
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
|
||||
/// CUDA contains its own atomic operations
|
||||
///
|
||||
template <typename T>
|
||||
class DeviceAdapterAtomicArrayImplementation<T, vtkm::cont::DeviceAdapterTagCuda>
|
||||
{
|
||||
public:
|
||||
VTKM_CONT
|
||||
DeviceAdapterAtomicArrayImplementation(
|
||||
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> handle)
|
||||
: Portal(handle.PrepareForInPlace(vtkm::cont::DeviceAdapterTagCuda()))
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_EXEC T Add(vtkm::Id index, const T& value) const
|
||||
{
|
||||
T* lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
|
||||
return this->vtkmAtomicAdd(lockedValue, value);
|
||||
}
|
||||
|
||||
VTKM_EXEC 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 this->vtkmCompareAndSwap(lockedValue, newValue, oldValue);
|
||||
}
|
||||
|
||||
private:
|
||||
using PortalType =
|
||||
typename vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>::template ExecutionTypes<
|
||||
vtkm::cont::DeviceAdapterTagCuda>::Portal;
|
||||
PortalType Portal;
|
||||
|
||||
__device__ vtkm::Int64 vtkmAtomicAdd(vtkm::Int64* address, const vtkm::Int64& value) const
|
||||
{
|
||||
return atomicAdd((unsigned long long*)address, (unsigned long long)value);
|
||||
}
|
||||
|
||||
__device__ vtkm::Int32 vtkmAtomicAdd(vtkm::Int32* address, const vtkm::Int32& value) const
|
||||
{
|
||||
return atomicAdd(address, value);
|
||||
}
|
||||
|
||||
__device__ vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
|
||||
const vtkm::Int32& newValue,
|
||||
const vtkm::Int32& oldValue) const
|
||||
{
|
||||
return atomicCAS(address, oldValue, newValue);
|
||||
}
|
||||
|
||||
__device__ vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
|
||||
const vtkm::Int64& newValue,
|
||||
const vtkm::Int64& oldValue) const
|
||||
{
|
||||
return atomicCAS((unsigned long long int*)address,
|
||||
(unsigned long long int)oldValue,
|
||||
(unsigned long long int)newValue);
|
||||
}
|
||||
};
|
||||
}
|
||||
} // end namespace vtkm::cont
|
||||
|
||||
#endif // vtk_m_cont_internal_DeviceAdapterAtomicArrayImplementationCuda_h
|
@ -30,6 +30,7 @@ set(headers
|
||||
ArrayTransfer.h
|
||||
ConnectivityExplicitInternals.h
|
||||
DeviceAdapterAlgorithmGeneral.h
|
||||
DeviceAdapterAtomicArrayImplementation.h
|
||||
DeviceAdapterDefaultSelection.h
|
||||
DeviceAdapterError.h
|
||||
DeviceAdapterListHelpers.h
|
||||
|
@ -27,6 +27,7 @@
|
||||
#include <vtkm/cont/ArrayHandleIndex.h>
|
||||
#include <vtkm/cont/ArrayHandleStreaming.h>
|
||||
#include <vtkm/cont/ArrayHandleZip.h>
|
||||
#include <vtkm/cont/internal/DeviceAdapterAtomicArrayImplementation.h>
|
||||
#include <vtkm/cont/internal/FunctorsGeneral.h>
|
||||
|
||||
#include <vtkm/exec/internal/ErrorMessageBuffer.h>
|
||||
@ -874,125 +875,8 @@ public:
|
||||
input, values_output, values_output);
|
||||
}
|
||||
};
|
||||
}
|
||||
}
|
||||
} // 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
|
||||
DeviceAdapterAtomicArrayImplementation(
|
||||
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> handle)
|
||||
: Iterators(IteratorsType(handle.PrepareForInPlace(DeviceTag())))
|
||||
{
|
||||
}
|
||||
|
||||
T Add(vtkm::Id index, const T& value) const
|
||||
{
|
||||
T* lockedValue;
|
||||
#if defined(_ITERATOR_DEBUG_LEVEL) && _ITERATOR_DEBUG_LEVEL > 0
|
||||
using IteratorType = typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType;
|
||||
typename IteratorType::pointer temp =
|
||||
&(*(Iterators.GetBegin() + static_cast<std::ptrdiff_t>(index)));
|
||||
lockedValue = temp;
|
||||
return vtkmAtomicAdd(lockedValue, value);
|
||||
#else
|
||||
lockedValue = (Iterators.GetBegin() + index);
|
||||
return vtkmAtomicAdd(lockedValue, value);
|
||||
#endif
|
||||
}
|
||||
|
||||
T CompareAndSwap(vtkm::Id index, const T& newValue, const T& oldValue) const
|
||||
{
|
||||
T* lockedValue;
|
||||
#if defined(_ITERATOR_DEBUG_LEVEL) && _ITERATOR_DEBUG_LEVEL > 0
|
||||
using IteratorType = typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType;
|
||||
typename IteratorType::pointer temp =
|
||||
&(*(Iterators.GetBegin() + static_cast<std::ptrdiff_t>(index)));
|
||||
lockedValue = temp;
|
||||
return vtkmCompareAndSwap(lockedValue, newValue, oldValue);
|
||||
#else
|
||||
lockedValue = (Iterators.GetBegin() + index);
|
||||
return vtkmCompareAndSwap(lockedValue, newValue, oldValue);
|
||||
#endif
|
||||
}
|
||||
|
||||
private:
|
||||
using PortalType =
|
||||
typename vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>::template ExecutionTypes<
|
||||
DeviceTag>::Portal;
|
||||
using IteratorsType = vtkm::cont::ArrayPortalToIterators<PortalType>;
|
||||
IteratorsType Iterators;
|
||||
|
||||
#if defined(VTKM_MSVC) //MSVC atomics
|
||||
vtkm::Int32 vtkmAtomicAdd(vtkm::Int32* address, const vtkm::Int32& value) const
|
||||
{
|
||||
return InterlockedExchangeAdd(reinterpret_cast<volatile long*>(address), value);
|
||||
}
|
||||
|
||||
vtkm::Int64 vtkmAtomicAdd(vtkm::Int64* address, const vtkm::Int64& value) const
|
||||
{
|
||||
return InterlockedExchangeAdd64(reinterpret_cast<volatile long long*>(address), value);
|
||||
}
|
||||
|
||||
vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
|
||||
const vtkm::Int32& newValue,
|
||||
const vtkm::Int32& oldValue) const
|
||||
{
|
||||
return InterlockedCompareExchange(
|
||||
reinterpret_cast<volatile long*>(address), newValue, oldValue);
|
||||
}
|
||||
|
||||
vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
|
||||
const vtkm::Int64& newValue,
|
||||
const vtkm::Int64& oldValue) const
|
||||
{
|
||||
return InterlockedCompareExchange64(
|
||||
reinterpret_cast<volatile long long*>(address), newValue, oldValue);
|
||||
}
|
||||
|
||||
#else //gcc built-in atomics
|
||||
|
||||
vtkm::Int32 vtkmAtomicAdd(vtkm::Int32* address, const vtkm::Int32& value) const
|
||||
{
|
||||
return __sync_fetch_and_add(address, value);
|
||||
}
|
||||
|
||||
vtkm::Int64 vtkmAtomicAdd(vtkm::Int64* address, const vtkm::Int64& value) const
|
||||
{
|
||||
return __sync_fetch_and_add(address, value);
|
||||
}
|
||||
|
||||
vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
|
||||
const vtkm::Int32& newValue,
|
||||
const vtkm::Int32& oldValue) const
|
||||
{
|
||||
return __sync_val_compare_and_swap(address, oldValue, newValue);
|
||||
}
|
||||
|
||||
vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
|
||||
const vtkm::Int64& newValue,
|
||||
const vtkm::Int64& oldValue) const
|
||||
{
|
||||
return __sync_val_compare_and_swap(address, oldValue, newValue);
|
||||
}
|
||||
|
||||
#endif
|
||||
};
|
||||
} // namespace internal
|
||||
|
||||
/// \brief Class providing a device-specific support for selecting the optimal
|
||||
/// Task type for a given worklet.
|
||||
|
152
vtkm/cont/internal/DeviceAdapterAtomicArrayImplementation.h
Normal file
152
vtkm/cont/internal/DeviceAdapterAtomicArrayImplementation.h
Normal file
@ -0,0 +1,152 @@
|
||||
//============================================================================
|
||||
// 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 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
|
||||
// Copyright 2018 UT-Battelle, LLC.
|
||||
// Copyright 2018 Los Alamos National Security.
|
||||
//
|
||||
// Under the terms of Contract DE-NA0003525 with NTESS,
|
||||
// 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_cont_internal_DeviceAdapterAtomicArrayImplementation_h
|
||||
#define vtk_m_cont_internal_DeviceAdapterAtomicArrayImplementation_h
|
||||
|
||||
#include <vtkm/cont/ArrayHandle.h>
|
||||
#include <vtkm/cont/ArrayPortalToIterators.h>
|
||||
#include <vtkm/cont/StorageBasic.h>
|
||||
|
||||
#include <vtkm/internal/Configure.h>
|
||||
#include <vtkm/internal/Windows.h>
|
||||
|
||||
#include <vtkm/Types.h>
|
||||
|
||||
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
|
||||
{
|
||||
using PortalType =
|
||||
typename vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>::template ExecutionTypes<
|
||||
DeviceTag>::Portal;
|
||||
using IteratorsType = vtkm::cont::ArrayPortalToIterators<PortalType>;
|
||||
IteratorsType Iterators;
|
||||
|
||||
public:
|
||||
VTKM_CONT
|
||||
DeviceAdapterAtomicArrayImplementation(
|
||||
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic> handle)
|
||||
: Iterators(IteratorsType(handle.PrepareForInPlace(DeviceTag())))
|
||||
{
|
||||
}
|
||||
|
||||
T Add(vtkm::Id index, const T& value) const
|
||||
{
|
||||
T* lockedValue;
|
||||
#if defined(_ITERATOR_DEBUG_LEVEL) && _ITERATOR_DEBUG_LEVEL > 0
|
||||
using IteratorType = typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType;
|
||||
typename IteratorType::pointer temp =
|
||||
&(*(Iterators.GetBegin() + static_cast<std::ptrdiff_t>(index)));
|
||||
lockedValue = temp;
|
||||
return this->vtkmAtomicAdd(lockedValue, value);
|
||||
#else
|
||||
lockedValue = (Iterators.GetBegin() + index);
|
||||
return this->vtkmAtomicAdd(lockedValue, value);
|
||||
#endif
|
||||
}
|
||||
|
||||
T CompareAndSwap(vtkm::Id index, const T& newValue, const T& oldValue) const
|
||||
{
|
||||
T* lockedValue;
|
||||
#if defined(_ITERATOR_DEBUG_LEVEL) && _ITERATOR_DEBUG_LEVEL > 0
|
||||
using IteratorType = typename vtkm::cont::ArrayPortalToIterators<PortalType>::IteratorType;
|
||||
typename IteratorType::pointer temp =
|
||||
&(*(Iterators.GetBegin() + static_cast<std::ptrdiff_t>(index)));
|
||||
lockedValue = temp;
|
||||
return this->vtkmCompareAndSwap(lockedValue, newValue, oldValue);
|
||||
#else
|
||||
lockedValue = (Iterators.GetBegin() + index);
|
||||
return this->vtkmCompareAndSwap(lockedValue, newValue, oldValue);
|
||||
#endif
|
||||
}
|
||||
|
||||
private:
|
||||
#if defined(VTKM_MSVC) //MSVC atomics
|
||||
vtkm::Int32 vtkmAtomicAdd(vtkm::Int32* address, const vtkm::Int32& value) const
|
||||
{
|
||||
return InterlockedExchangeAdd(reinterpret_cast<volatile long*>(address), value);
|
||||
}
|
||||
|
||||
vtkm::Int64 vtkmAtomicAdd(vtkm::Int64* address, const vtkm::Int64& value) const
|
||||
{
|
||||
return InterlockedExchangeAdd64(reinterpret_cast<volatile long long*>(address), value);
|
||||
}
|
||||
|
||||
vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
|
||||
const vtkm::Int32& newValue,
|
||||
const vtkm::Int32& oldValue) const
|
||||
{
|
||||
return InterlockedCompareExchange(
|
||||
reinterpret_cast<volatile long*>(address), newValue, oldValue);
|
||||
}
|
||||
|
||||
vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
|
||||
const vtkm::Int64& newValue,
|
||||
const vtkm::Int64& oldValue) const
|
||||
{
|
||||
return InterlockedCompareExchange64(
|
||||
reinterpret_cast<volatile long long*>(address), newValue, oldValue);
|
||||
}
|
||||
|
||||
#else //gcc built-in atomics
|
||||
|
||||
vtkm::Int32 vtkmAtomicAdd(vtkm::Int32* address, const vtkm::Int32& value) const
|
||||
{
|
||||
return __sync_fetch_and_add(address, value);
|
||||
}
|
||||
|
||||
vtkm::Int64 vtkmAtomicAdd(vtkm::Int64* address, const vtkm::Int64& value) const
|
||||
{
|
||||
return __sync_fetch_and_add(address, value);
|
||||
}
|
||||
|
||||
vtkm::Int32 vtkmCompareAndSwap(vtkm::Int32* address,
|
||||
const vtkm::Int32& newValue,
|
||||
const vtkm::Int32& oldValue) const
|
||||
{
|
||||
return __sync_val_compare_and_swap(address, oldValue, newValue);
|
||||
}
|
||||
|
||||
vtkm::Int64 vtkmCompareAndSwap(vtkm::Int64* address,
|
||||
const vtkm::Int64& newValue,
|
||||
const vtkm::Int64& oldValue) const
|
||||
{
|
||||
return __sync_val_compare_and_swap(address, oldValue, newValue);
|
||||
}
|
||||
|
||||
#endif
|
||||
};
|
||||
}
|
||||
} // end namespace vtkm::cont
|
||||
|
||||
#endif // vtk_m_cont_internal_DeviceAdapterAtomicArrayImplementation_h
|
Loading…
Reference in New Issue
Block a user