The `UnitTestVTKDataSetReader` test reads in .vtk files of multiple
formats. The origins of this test predate the use of git-lfs and a data
directory, so the test defined in the source code several strings
containing the contents of the files.
That is no longer necessary. This change removes those strings from the
file and creates actual files in the data directory. This is easier to
manage as well as more convenient if you want to use the same data in a
different test or just take a look at the data in a different program
(such as VisIt or ParaView).
The testing helper class provided a method named `GetTestDataBasePath`
that returned the base path to all the data files stored in the VTK-m
repo. This is fine, but it was a little cumbersome to build filenames.
To make things easier, there is now a new method named `DataPath` that
takes a string of the filename (or, rather, subpath) to the file in that
directory and automatically builds the path to it.
4345fe26b Store the number of bits of a BitField in the Buffer's metadata
da0403be7 Add metadata to Buffer object.
a84891cd3 Update ArrayHandleBitField to new array style with Buffer
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2218
The number of bits in a `BitField` cannot be directly implied from the
size of the buffer (because the buffer gets padded to the nearest sized
word). Thus, the `BitField stored the number of bits in its own
internals.
Unfortunately, that caused issues when passing the `BitField` data
between it and an `ArrayHandleBitField`. If the `ArrayHandleBitField`
resized itself, the `BitField` would not see the new size because it
ignored the new buffer size.
To get around this problem, `BitField` now declares its own
`BufferMetaData` that stores the number of bits. Now, since the number
of bits is stored in the `Buffer` object, it is sufficient to just share
the `Buffer` to synchronize all of the state.
One of the goals of the `Buffer` object is to allow sharing of data
among objects that will interpret the data differently or give a
different interface over the data. However, when sharing only the array,
important metadata can become lost.
Provide a field that can store some custom metadata in the buffer object
so that the rest of the state can follow the buffer object around.
6cbcb9f5d Fix behavior of Cuda AtomicLoad with SequentiallyConsistent
7573d4ed5 Fix compiler warnings
147dd24d0 Remove ARM intrinsics in MSVC
2229c22f4 Avoid invalid Kokkos atomic calls
3b147878f Always use our implementation of Cuda atomics
9e6fe8fb6 Add memory order semantics to atomic functions
d2ac4b860 Be more careful in casting with Atomic functions
13056b3af Deprecate AtomicInterfaceControl and AtomicInterfaceExecution
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2223
According to Allie Vacanti, a sequentially consistent load requires a
full fence on Cuda hardware to be conforming.
Also improved the documentation of `MemoryOrder` based on Allie's
suggestion.
Also removed the `Consume` memory order based on Allie's suggestion. It
is tricky to use correctly, and most implementations just regress to the
safer `Acquired` behavior anyway.
Previously, if Kokkos was enabled we always used Kokkos atomics.
However, if a user, for some reason, compiled with a version of Kokkos
that was _not_ compiled for Cuda and we turned on our own Cuda device
adapter, that would cause a problem. The old code assumed Kokkos would
create the Cuda version of the atomics, but it would not. Thus, there
would be no atomics for Cuda.
Resolved this problem by switching the order in which we try to define
atomics. Use our version of atomics whenever compiling for a Cuda
device. Otherwise, try to compile the Kokkos version (and if that is not
available, use ours as before.
To ensure correctness of an atomic function, it is often necessary to
force a compiler to order operations correctly. However, doing so may
slow things down, so it is helpful to relax these constraints if they
are not necessary. Add the ability to select the correct memory order
constraints in the atomic functions.
Previously, the atomic functions accepted any type as its operand and
then cast that to the storage type. That could cause some rather
unexpected behavior. For example, casting a floating point number to an
integer might not give you the behavior as expected. So that behavior as
been removed and now the operand has to match the pointer.
However, all the currently supported atomics are unsigned, and there are
many reasons that it might be easier to use signed as operands. (For
example, C literals are signed.) Thus, a second condition that allows
the sign to be swapped has been added so that you don't get annoying
signed/unsigned conversion warnings.
Now that we have the functions in `vtkm/Atomic.h`, we can deprecate (and
eventually remove) the more cumbersome classes `AtomicInterfaceControl`
and `AtomicInterfaceExecution`.
Also reversed the order of the `expected` and `desired` parameters of
`vtkm::AtomicCompareAndSwap`. I think the former order makes more sense
and matches more other implementations (such as `std::atomic` and the
GCC `__atomic` built ins). However, there are still some non-deprecated
classes with similar methods that cannot easily be switched. Thus, it's
better to be inconsistent with most other libraries and consistent with
ourself than to be inconsitent with ourself.
If the Kokkos device is enabled, we use the Kokkos atomic functions for
implementation since Kokkos has already ported all of these to each
device.
If Kokkos is compiled with CUDA, then the functions are marked with
`__device__` and `__host__`. Makes sense right?
Unless we are trying to use the atomic functions on host code compiled
only for the host. In that case, modifiers like `__device__` will cause
compiler errors.
So we want to disable Kokkos from using `__device__`, but only if CUDA
is enabled and we are not using the CUDA compiler. Had to hack things up
to get that to work.
Now that we have atomic free functions (e.g. `vtkm::AtomicAdd()`), we no
longer need special implementations for control and each execution
device. (Well, technically we do have special implementations for each,
but they are handled with compiler directives in the free functions.)
Convert the old atomic interface classes (`AtomicInterfaceControl` and
`AtomicInterfaceExecution`) to use the new atomic free functions. This
will allow us to test the new atomic functions everywhere that atomics
are used in VTK-m.
Once verified, we can deprecate the old atomic interface classes.
Previously, all atomic functions were stored in classes named
`AtomicInterfaceControl` and `AtomicInterfaceExecution`, which required
you to know at compile time which device was using the methods. That in
turn means that anything using an atomic needed to be templated on the
device it is running on.
That can be a big hassle (and is problematic for some code structure).
Instead, these methods are moved to free functions in the `vtkm`
namespace. These functions operate like those in `Math.h`. Using
compiler directives, an appropriate version of the function is compiled
for the current device the compiler is using.
There is a limitation in Windows builds using VS2019 where libraries cannot be
bigger than 4GiB. This is normally not an issue but in `VTKm` due to its strong
template usage libraries can reach that size.
The `VTKm` filter library is can easily reach that size and it will halt the
build
This MR tries to avoid reaching those sizes for now by splitting the filter
library into four smaller libraries.
The proposal scheme is:
It splits vtkm-filter into:
- vtkm-common, Classes that are dependencies of other filter libs.
- vtkm-contour, Contour class and its instantiations.
- vtkm-contour, Gradient class and its instantiations.
- vtkm-extra, Classes other than Contour or Gradient that are
not dependencies.
Signed-off-by: Vicente Adolfo Bolea Sanchez <vicente.bolea@kitware.com>
ed41874cc Consolidate tests for base vtkm code that is device-specific
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2219
Some of the code in the base `vtkm` namespace is device specific. For
example, the functions in `Math.h` are customized for specific devices.
Thus, we want this code to be specially compiled and run on these
devices.
Previously, we made a header file and then added separate tests to each
device package. That was created before we had ways of running on any
device. Now, it is much easier to compile the test a single time for all
devices and use the `ALL_BACKENDS` feature of `vtkm_unit_tests` CMake
function to automatically create the test for all devices.
8983154e9 Add DeepCopy to ArrayHandle
694ba7e92 Change ArrayCopy to deep copy Buffer objects where possible
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2212
2d1b609b3 Use Ubuntu instead of rhel8 for cuda+kokkos
769248583 Make sure we use c++14 when using CUDA 11+
64efa6401 Kokkos: make sure we don't pass multiple rdc flags
b2f4c8e5e Switch -O3 to -O2 on Linux with Cuda 10
db57ed26a Fix warnings
452f61e29 Add Kokkos backend
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2164
`ArrayHandle::DeepCopy` creates a new `ArrayHandle` of the same type and
deep copies the data into it.
This functionality is similar to `ArrayCopy`. However, it can be used
without having to compile for the device on which the copy happens.
Now that the data in an `ArrayHandle` is stored in `Buffer` objects, we
now have a more efficient way of doing deep copies of memory. Rather
than call `Algorithm::Copy`, which iterates over the array and copies
each item, `ArrayCopy` now uses the `Buffer` interface to do direct
device-to-device (or host-to-host) mem copies. This should be more
efficent and take less time to compile.
Note that this direct `Buffer` copy only works if the two `ArrayHandle`s
are of the same type. If they are different, `ArrayCopy` still has to
fall back to using `Algorithm::Copy`.
Also note that not all `ArrayHandle`s are using the new `ArrayHandle`
interface (and therefore not using `Buffer` objects). Thus, a fallback
is still available for old `ArrayHandle` types.
This has no real change in the operation, but it will simplify code as
we convert `ArrayHandle`s to the new type. We will be able to write
simple runtime code rather than complex metaprogramming to determine the
number of buffers to use.
For `make_ArrayHandle` and `make_Field` when it is determined that the
data can be safely moved, just silently move instead of copy instead of
printing a log message saying the copy flag will be ignored.
Also fix an issue with `make_ArrayHandle` when the data was not moved
when it could have been.
The recent version of ArrayHandleBasic allocates typeless arrays without
any initialization. This can cause issues with types that require a
constructor. The UnitTestVariantArrayHandle was trying to create an
ArrayHandle with an std::string, and the uninitialized strings were
causing crashes on some platforms.
We have made several improvements to adding data into an `ArrayHandle`.
## Moving data from an `std::vector`
For numerous reasons, it is convenient to define data in a `std::vector`
and then wrap that into an `ArrayHandle`. It is often the case that an
`std::vector` is filled and then becomes unused once it is converted to an
`ArrayHandle`. In this case, what we really want is to pass the data off to
the `ArrayHandle` so that the `ArrayHandle` is now managing the data and
not the `std::vector`.
C++11 has a mechanism to do this: move semantics. You can now pass
variables to functions as an "rvalue" (right-hand value). When something is
passed as an rvalue, it can pull state out of that variable and move it
somewhere else. `std::vector` implements this movement so that an rvalue
can be moved to another `std::vector` without actually copying the data.
`make_ArrayHandle` now also takes advantage of this feature to move rvalue
`std::vector`s.
There is a special form of `make_ArrayHandle` named `make_ArrayHandleMove`
that takes an rvalue. There is also a special overload of
`make_ArrayHandle` itself that handles an rvalue `vector`. (However, using
the explicit move version is better if you want to make sure the data is
actually moved.)
## Make `ArrayHandle` from initalizer list
A common use case for using `std::vector` (particularly in our unit tests)
is to quickly add an initalizer list into an `ArrayHandle`. Now you can
by simply passing an initializer list to `make_ArrayHandle`.
## Deprecated `make_ArrayHandle` with default shallow copy
For historical reasons, passing an `std::vector` or a pointer to
`make_ArrayHandle` does a shallow copy (i.e. `CopyFlag` defaults to `Off`).
Although more efficient, this mode is inherintly unsafe, and making it the
default is asking for trouble.
To combat this, calling `make_ArrayHandle` without a copy flag is
deprecated. In this way, if you wish to do the faster but more unsafe
creation of an `ArrayHandle` you should explicitly express that.
This requried quite a few changes through the VTK-m source (particularly in
the tests).
## Similar changes to `Field`
`vtkm::cont::Field` has a `make_Field` helper function that is similar to
`make_ArrayHandle`. It also features the ability to create fields from
`std::vector`s and C arrays. It also likewise had the same unsafe behavior
by default of not copying from the source of the arrays.
That behavior has similarly been depreciated. You now have to specify a
copy flag.
The ability to construct a `Field` from an initializer list of values has
also been added.
While compiling UnitTestVariantArrayHandle, some versions of gcc
(between 6 and 8, I think) gave a warning like the following:
```
../vtkm/cont/StorageVirtual.h:227:12: warning: 'vtkm::Id vtkm::cont::internal::detail::StorageVirtualImpl<T, S>::GetNumberOfValues() const [with T = std::__cxx11::basic_string<char>; S = vtkm::cont::StorageTagImplicit<{anonymous}::UnusualPortal<std::__cxx11::basic_string<char> > >]' declared 'static' but never defined [-Wunused-function]
```
This warning makes no sense because it is refering to a method that is
not declared static. (In fact, it overrides a virtual method.)
I believe this is an obscure bug in these versions of gcc. I found a
[stackoverflow post] that seems to have the same problem, but no
workaround was found.
The warning originated from code that had little effect. It was part of
a test with a custom ArrayHandle storage type that was already disabled
for other reasons. Just removed the code.
[stackoverflow post]: https://stackoverflow.com/questions/56615695/how-to-fix-declared-static-but-never-defined-on-member-function
c689a68c5 Suppress bad deprecation warnings in MSVC
a3f23a03b Do not cast to ArrayHandleVirtual in VariantArrayHandle::CastAndCall
f6b13df51 Support coordinates of both float32 and float64
453e31404 Deprecate ArrayHandleVirtualCoordinates
be7f06bbe CoordinateSystem data is VariantArrayHandle
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2177
The Microsoft compiler has this annoying and stupid behavior where if
you have a generic templated method/function and that method is
instantiated with a deprecated class, then the compiler will issue a
C4996 warning even if the calling code is suppressing that warning
(because, for example, you are implementing other deprecated code and
the use is correct). There is no way around this other than suppressing
the warnings for all uses of the templated method.
A recent change to the continuous integration setup has caused
`UnitTestArrayPortalValueReference` to fail on one platform.
The problem was that the test was doing two unsafe things with the
right-shift assignment operator. The first unsafe thing was using itself
as the operand.
```cpp
ref >>= ref;
```
This causes clang to give a "self assign overload" warning. Using a
variable as its own operand for a compute assign operation isn't great
style, but for some operations it can cause weird errors. The reason for
the warning is that for a true integer shift operation, the compiler
will recognize that the result should be 0. So, when using this on base
integer types, you will get 0. But overloads can give you something
different, so that might lead to unexpected results. Because we _are_
using an overload (for the `ArrayPortalValueReference` class), the
compiler tells us we can get potentially unexpected results.
OK. That seems like something we can safely ignore (and were ignoring
for some time). After all, the whole point of the
`ArrayPortalValueReference` operators is to behave exactly the same as
the values they wrap. That brings us to the second unsafe thing the test
was doing: using an invalid value as the right hand operation. And this
is where things get weird.
The test was specifically failing when being run on `Int32`. It was
setting the underlying value for `ref` to be `1000` to start with. So
the expression `ref >>= ref` was trying to right shift `ref` by 1000
bits. Logically, this should of course give you 0. However, shifting a
number by more bits than exist causes undefined behavior (c.f.
https://wiki.sei.cmu.edu/confluence/display/c/INT34-C.+Do+not+shift+an+expression+by+a+negative+number+of+bits+or+by+greater+than+or+equal+to+the+number+of+bits+that+exist+in+the+operand).
You might not get back the expected value, and this is exactly what was
happening.
What I think happened was that the compiler determined that any valid
shift would be contained in 5 bits, so it truncated the value (1000) to
the least signifcant 5 bits (which happens to be 8). It then shifted
1000 by 8 and got the value 3 instead of 0.
The fix picks an operand number that is sure to be valid.
We are moving to deprecate `ArrayHandleVirtual`, so we are removing the
feature where `VariantArrayHandle::CastAndCall` automatically casts to
an `ArrayHandleVirtual` if possible.
The big reason to make this change now (as opposed to later when
`ArrayHandleVirtual` is deprecated) is to improve compile times.
This prevents us from having to compile an extra code path using
`ArrayHandleVirtual`.
`CoordinateSystem` differed from `Field` in that its `GetData`
method returned an `ArrayHandleVirtualCoordinates` instead of
a `VariantArrayHandle`. This is probably confusing since
`CoordianteSystem` inherits `Field` and has a pretty dramatic
difference in this behavior.
In preparation to deprecate `ArrayHandleVirtualCoordinates`, this
changes `CoordiantSystem` to be much more like `Field`. (In the
future, we may change the `CoordinateSystem` to point to a `Field`
rather than be a special `Field`.)
A method named `GetDataAsMultiplexer` has been added to
`CoordinateSystem`. This method allows you to get data from
`CoordinateSystem` as a single array type without worrying
about creating functors to handle different types and without
needing virtual methods.
18b5be92d Fix issue with CUDA and ArrayHandleMultiplexer
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2168
As a programming convenience, all `vtkm::cont::DataSet` written by
`vtkm::io::VTKDataSetWriter` were written as a structured grid. Although
technically correct, it changed the structure of the data. This meant that
if you wanted to capture data to run elsewhere, it would run as a different
data type. This was particularly frustrating if the data of that structure
was causing problems and you wanted to debug it.
Now, `VTKDataSetWriter` checks the type of the `CoordinateSystem` to
determine whether the data should be written out as `STRUCTURED_POINTS`
(i.e. a uniform grid), `RECTILINEAR_GRID`, or `STRUCTURED_GRID`
(curvilinear).
`ArrayHandle::PrepareForOutput` often has to reallocate the array to the
specified size. Previously, this allocation was not happening with the
`Token` that is passed to `PrepareForOutput`. If the `ArrayHandle` is
already attached or enqueued for that `Token`, then the allocation would
deadlock.
You can now pass a `Token` object to `Allocate`, which is what
`PrepareForOutput` does.
When you try to call the `Reduce` operation in the CUDA device adapter
with a sufficently complex interator type, you get a compile error
that says `error: cannot pass an argument with a user-provided
copy-constructor to a device-side kernel launch`.
This appears to be a bug in either nvcc or Thrust. I believe it is
related to the following reported issues:
* https://github.com/thrust/thrust/issues/928
* https://github.com/thrust/thrust/issues/1044
Work around this problem by making a special condition for calling
`Reduce` with an `ArrayHandleMultiplexer` that calls the generic
algorithm in `DeviceAdapterAlgorithmGeneral` instead of the algorithm in
Thrust.
The implementation of LagrangianFilter uses some `ArrayHandle`s
declared as `static` in the .hxx header file. One of the bad
consequences of this is that there is no control over when the
arrays are freed. We have seen where these arrays get freed
after the CUDA system is closed, which causes nasty things to
happen as the program closes.
This works around the problem until a fix is implemnted.
Often when a user gives memory to an `ArrayHandle`, she wants data to be
written into the memory given to be used elsewhere. Previously, the
`Buffer` objects would delete the given buffer as soon as a write buffer
was created elsewhere. That was a problem if a user wants VTK-m to write
results right into a given buffer.
Instead, when a user provides memory, "pin" that memory so that the
`ArrayHandle` never deletes it.
The buffer class encapsulates the movement of raw C arrays between
host and devices.
The `Buffer` class itself is not associated with any device. Instead,
`Buffer` is used in conjunction with a new templated class named
`DeviceAdapterMemoryManager` that can allocate data on a given
device and transfer data as necessary. `DeviceAdapterMemoryManager`
will eventually replace the more complicated device adapter classes
that manage data on a device.
The code in `DeviceAdapterMemoryManager` is actually enclosed in
virtual methods. This allows us to limit the number of classes that
need to be compiled for a device. Rather, the implementation of
`DeviceAdapterMemoryManager` is compiled once with whatever compiler
is necessary, and then the `RuntimeDeviceInformation` is used to
get the correct object instance.
143e3d39a remove unused type alias
01a448663 Merge branch 'master' into uniform_real
c67e5bb12 fixe warnings about implicit type conversion
1e4294392 Add deterministic seed to avoid potential spurious failure
5b0e309b9 the random source is still 64 bits
cc3061bab Avoid calling ReadPortal() all the time
9bf6dea22 remove inline initialization of seed
e69308047 Add statistics base testing, add Flot32 RNG
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2148
If a test throws any unexpected exception, the test is supposed to
detect that and fail. For the STL exceptions, the test failed to return
an error code. Fix that.
For some reason, `UnitTestContourTreeUniformAugmentedFilterCUDA` was
deadlocking on some of the dashboards when `VTKM_ASSERT` was changed to
be empty when compiling CUDA kernels (for compiler performance reasons).
Fix this by calling `assert` directly in this case.
For the life of me, I cannot figure out why this would be an issue.
Clearly the `assert` is never actually called or else the test would
error out (unless a special condition in CUDA is causing it to be
hidden). But if you take out the code, something changes to lock up the
kernel.
`assert` is supported on recent CUDA cards, but compiling it appears to be
very slow. By default, the `VTKM_ASSERT` macro has been disabled whenever
compiling for a CUDA device (i.e. when `__CUDA_ARCH__` is defined).
Asserts for CUDA devices can be turned back on by turning the
`VTKm_NO_ASSERT_CUDA` CMake variable off. Turning this CMake variable off
will enable assertions in CUDA kernels unless there is another reason
turning off all asserts (such as a release build).
The VS compiler gives a warning when you override a deprecated method.
Changed the StartScene and EndScene methods in Mapper to be non-virtual
with empty implementations. Deleted the corresponding methods from all
subclasses.
Many of the operations of `VariantArrayHandleBase` are not dependent on
the TypeList parameter of the class. Still others can operate just as
well by providing a type list to a method. Thus, it is convenient to
create a superclass that is not templated. That allows us to pass around
a `VariantArrayHandle` when the type list does not matter.
This superclass is called `VariantArrayHandleCommon` because "base" was
already taken.
This makes it clear that it returns true for an invalid array handle.
The previous name implied that it was looking for an ArrayHandle in some
"valid" set, which is the opposite.
0a4317709 Fix issues of calling __host__ from __device__
dd4d88cd5 Fix warnings about comparing floating point values
cd4b59059 Fix warnings about data type conversion
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2146
df0881631 minor coding style improvment
68c18f154 remove redundant #include
ba35520b1 use condition number for error tolerance
ee422b7f2 move data member initialization to constructors
e0c4db209 add VTKM_EXEC_CONT to constructor
c25f0b4bb Rename StatisticalMoments. more unit tests
5eeeb4791 Change StatState's API to better conform to VTKm's coding standard
794db28ad Merge branch 'master' into stats_by_key
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2128
The ArrayPortalMultiplexer always called Set on the delegate portals.
However, sometimes the portals were not supported. Instead, only call
when the delegate is supported.
Probably the "right" thing to do would be to check all posible portals
for support, but that sounds like it would slow down the compile.
Previously, the VariantArrayHandle::AsMultiplexer operation silently
failed and returned an invalid ArrayHandleMultiplexer. This is now
changed to throw an exception if the desired ArrayHandleMultiplexer
cannot hold the VariantArrayHandle's array.
This merge request is Phase 1 of several to implement the distributed parallel
contour tree in VTKm. This merge requests adds the base outline for the
algorithm. The implementation of the details of the algorithm in the
BoundaryRestrictedAugmentedContourTree.h is currently still missing.
However, these will require a substantial (~3000) lines of additional code.
The goal is to stage the integration process across merge requests to make
the review process simpler.
de3bda373 Use deque instead of list for ArrayHandle queue
498d44548 Pass Token::Reference by value
c32c9e8e8 Fix deadlock when changing device during read
99e14ab8a Add proper enqueuing of Tokens for ArrayHandle
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2130
16e04db15 FlyingEdges: Handle when pass 1 generates no intersections
8ebea33c7 FlyingEdges: Handle when a dimensions size is <= 3
60eff58bd Correct warning in TransferToOpenGL found by the examples
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2125
c0dee7402 make it explicit that we are using 64-bit unsigned integer in bit op
34f350588 Added changelog
e9f584a91 ArrayHandleRandomUniformReal
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !2116
934f085e0 Build diy as a library
f0a37ac6a Merge branch 'upstream-diy' into diy-mpi-nompi
7687aabf8 diy 2020-06-05 (b62915aa)
6ca2b9f87 Point to new version of Diy
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2123
Because ArrayHandle currently only supports one device at a time, it is
possible that a `PrepareForInput` might actually need to wait for write
access so that it could move the data between devices. However, we don't
want the `Token` to be attached for writing because that could block
other read operations.
To get around this, add a hack to WaitToWrite to allow it to attach for
reading instead of writing.
An issue that was identified for the thread safety of `ArrayHandle` is
that if several threads are waiting to use an `ArrayHandle`, there might
be an expectation of the order in which the operations happen. For
example, if one thread is modifying the contents of an `ArrayHandle` and
another is reading those results, we would need the first one to start
before the second one.
To solve this, a queue is added to `ArrayHandle` such that when waiting
to read or write an `ArrayHandle` the `Token` has to be at the top of
the queue in addition to other requirements being met.
Additionally, an `Enqueue` method is added to add a `Token` to the queue
without blocking. This allows a control thread to queue the access and
then spawn a thread where the actual work will be done. As long as
everything is enqueued on the main thread, the operations will happen in
the expected order.
ce3a64b41 fix comparison
ab6bb14d3 Merge branch 'respect_max_distance' of gitlab.kitware.com:mclarsen/vtk-m into respect_max_distance
a6ed41027 use bounds based epsilon for tiny meshes
59775c4eb respect ray max distance
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Sudhanshu Sane <ssane@cs.uoregon.edu>
Merge-request: !2129
Initially, the probe filter would simply not set a value if a sample was
outside the input `DataSet`. This is not great as the memory could be
left uninitalized and lead to unpredictable results. The testing
compared these invalid results to 0, which seemed to work but is
probably unstable.
This was partially fixed by a previous change that consolidated to
mapping of cell data with a general routine that permuted data. However,
the fix did not extend to point data in the input, and it was not
possible to specify a particular invalid value.
This change specifically updates the probe filter so that invalid values
are set to a user-specified value.
