`ArrayPortalToIterators` broken on MSVC + CUDA
When compiling with MSVC, ArrayPortalToIterators
uses a platform-specific checked-iterator class to wrap pointer accesses:
// The MSVC compiler issues warnings when using raw pointer math when in
// debug mode. To keep the compiler happy (and add some safety checks),
// wrap the iterator in checked_array_iterator.
using IteratorType = stdext::checked_array_iterator<IterType>;
This is incompatible with CUDA, since the checked_array_iterator
lacks __host__ __device__
markup:
src\vtk-m\vtkm/Algorithms.h(143): warning #2979-D: calling a __host__ function from a __host__ __device__ function is not allowed
detected during:
instantiation of "IterT vtkm::UpperBound(IterT, IterT, const T &, Comp) [with IterT=stdext::checked_array_iterator<const vtkm::Int64 *>, T=vtkm::Id, Comp=vtkm::SortLess]"
We have two options here:
- Remove the checked iterator and just store the pointer on all platforms.
- Only enable the checked iterator when compiling without
nvcc
.
I'm leaning towards 1. The implementation will be cleaner, and memcheck
/cuda-memcheck
are great at finding the sorts of issues the checked_array_iterator
detects. IIRC, the warnings mentioned in the comment are disabled via MSVC's SECURE_NO_WARNINGS
flags.
Option 2 seems to get into some weird binary compatibility issues, and prevents iterators from being able to be passed between host <-> device.
Does anyone else have input on this?