0b2ddb83b UnitTestBounds custom compile flags expressed via generator expression
1ed5dfca0 vtkm/Math frexp(float) version works now with HIP
ff381bf8b vtkm/Swap works with hip
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !2301
0b1c48a3e Identify number of partitions in filter execution
b33c54bf6 Add ScheduleTask to performance log
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2302
When you execute a filter, the default behavior is to do the execution
on each partition of the data set independently. This code path is
followed even for non-partitioned data; the `DataSet` is wrapped in a
`PartitionedDataSet` of one partition.
Make performance logging a bit more clear by only giving one scoped log
for a basic `DataSet` and recording the number of partitions executed in
the log.
When `DeviceAdapterAlgorithm::ScheduleTask` was called directly (i.e.
not through `Schedule`), nothing was added to the log. Adding
`VTKM_LOG_SCOPE` to these methods so that all scheduling is added to the
performance log.
f66dc780c Update ArrayHandleImplicit to new array style with Buffer
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2246
- It fixes a lingering error triggered with BenchContour
- It reenables BenchContour
Signed-off-by: Vicente Adolfo Bolea Sanchez <vicente.bolea@kitware.com>
The old atomic compare and swap operations (`vtkm::AtomicCompareAndSwap`
and `vtkm::exec::AtomicArrayExecutionObject::CompareAndSwap`) had an
order of arguments that was confusing. The order of the arguments was
shared pointer (or index), desired value, expected value. Most people
probably assume expected value comes before desired value. And this
order conflicts with the order in the `std` methods, GCC atomics, and
Kokkos.
Change the interface of atomic operations to be patterned off the
`std::atomic_compare_exchange` and `std::atomic<T>::compare_exchange`
methods. First, these methods have a more intuitive order of parameters
(shared pointer, expected, desired). Second, rather than take a value
for the expected and return the actual old value, they take a pointer to
the expected value (or reference in `AtomicArrayExecutionObject`) and
modify this value in the case that it does not match the actual value.
This makes it harder to mix up the expected and desired parameters.
Also, because the methods return a bool indicating whether the value was
changed, there is an additional benefit that compare-exchange loops are
implemented easier.
For example, consider you want to apply the function `MyOp` on a
`sharedValue` atomically. With the old interface, you would have to do
something like this.
```cpp
T oldValue;
T newValue;
do
{
oldValue = *sharedValue;
newValue = MyOp(oldValue);
} while (vtkm::AtomicCompareAndSwap(sharedValue, newValue, oldValue) != oldValue);
```
With the new interface, this is simplfied to this.
```cpp
T oldValue = *sharedValue;
while (!vtkm::AtomicCompareExchange(sharedValue, &oldValue, MyOp(oldValue));
```
The `ReleaseResourcesExecution` method makes changes by booting data off
of the execution environment. But logically the array does not change.
It remains the same size with the same contents. Thus, it makes sense
for this to be a const method.
Also modified some deprecated methods a bit to remove some unnecessary
templates.
The `UnitTestParticleAdvectionFilter` was only testing advection in
structured cell sets (with either uniform or rectilinear points). This
did not test all paths (i.e. grid evaluators and locators) of particle
advection. Changed the test to also try grids with explicit cells.
Note that the explicit cells are created by just converting the
rectilinear grids to explicit grids. It would be better to get at least
one small dataset that starts as unstructured.
f5c5b6188 Log the computation of ranges in fields and arrays
bef55f8ee Log the conversion of num components to offsets
420b9d397 Log the building of arrays for scatters and masks
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !2283
c26e193fe Use kokkos provided routines for sort and copy
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2277
Previously vtk-m allowed users to issue atomic loads on constant
values which is problematic for the following reasons:
- can be a source of undefined behavior
- not supported by kokkos
This issue was detected when using kokkos HIP atomic implementation
When you invoke a worklet, a scoped performance log is automatically
generated. The previous version gave the name of the worklet, which is
technically everything you need to know. However, it is also convenient
to know what type of worklet it is (e.g. a field map worklet or a
topology map worklet). This can be determined by looking at the
definition of the worklet in the source code, but that is not practical
if, for example, you want to perform automated analysis.
This change prints out the type of the dispatcher instead of the
worklet. The full type of the dispatcher includes the worklet as a
template parameter, so you get both the dispatcher/worklet type and the
worklet itself. So previously you would get log lines like
```
Invoking Worklet: 'vtkm::worklet::vtkm::worklet::ExternalFaces::NumFacesPerCell'
```
They now look like this:
```
Invoking Worklet: 'vtkm::worklet::DispatcherMapTopology<vtkm::worklet::ExternalFaces::NumFacesPerCell>'
```
268a07bb0 remove test on large graph and sample data
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !2278
- It also remove termination on bench error
- It disables passing fields in Countour to skip an error.
Signed-off-by: Vicente Adolfo Bolea Sanchez <vicente.bolea@kitware.com>
932f8b82f Merge branch 'concurrent_union_find' of gitlab.kitware.com:ollielo/vtk-m into concurrent_union_find
eaff27195 minor changes on comments on findRoot()
9fef540ba minor changes on comments on findRoot()
916661e33 Merge branch 'concurrent_union_find' of gitlab.kitware.com:ollielo/vtk-m into concurrent_union_find
a036b0f98 Minor change based on code review.
9e68022ac Minor change based on code review.
1185f70bc Data Race Resolution
aed7e02da Merge branch 'concurrent_union_find' of gitlab.kitware.com:ollielo/vtk-m into concurrent_union_find
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !2248
Although `vtkm::internal::Variant` respected the trivially copyable
attribute of the types it contains, it was never totally trivial (i.e.
`std::is_trivial<Variant<...>>` was never true). The reason was that
`Variant` was initializing its `Index` parameter to signify that it was
not initialized. However, the fact that `Index` was initialized meant
that it was not trivially constructed.
Now, `Variant` type checks its types to see if they are all trivially
constructible. If so, it makes itself trivially constructible.
This means that `Index` may or may not be valid if `Variant` is
constructed without an argument. This in turn means that the result of
`Variant::IsValid` becomes undefined. That should be OK in practice.
`Index` will "point" to an uninitialized object, but that object is
trivially constructed anyway. However, that could cause problems if
developers used `IsValid` to determine if something is selected.
The new name reflects better what the underlying algorithm does. It also
helps prevent confusion about what types of data the locator is good
for. The old name suggested it only worked for structured grids, which
is not the case.
We would really like to be able to include `vtkm::cont::ColorTable` in
such a way that you don't have to compile device code (unless you are
actually compiling functions for the device). Thus, the `Map` functions
of `ColorTable` were in a special `ColorTable.hxx` that contains the
"implementation" for `ColorTable`.
That is confusing to many users. It is more clear to simply have `.h`
headers that do a specific thing. To achieve these two goals, the `Map`
functionality of `ColorTable` is separated out into its own header file.
So you don't need to be using a device compiler just to use `ColorTable`
(including `ColorTable.h`), but you do need to use a device compiler if
mapping values to colors (including `ColorTableMap.h`).
6a3ba4291 Fix warning about unused function
c6a4f9b79 Fix warning about return value
9465d2611 Adjust TransferFunction worklet to get color table as ExecObj argument
11996f133 Remove virtual methods from ColorTable
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2261
Previously, the `TransferFunction` worklet required you to call
`PrepareForExecution` on the `ColorTable` and give that to the
constructor. Bad things can happen if the `ColorTable` used in the
constructor does not match the device the worklet actually gets invoked
on. Change this to make the `ColorTable` a worklet argument of type
`ExecObj`. The dispatcher will automatically call
`ColorTable::PrepareForInput` and guarantee that the devices match.
Virtual methods are being deprecated, so remove their use from the
ColorTable classes. Instead of using a virtual method to look up a value
in the ColorTable, we essentially use a switch statement. This change
also simplified the code quite a bit.
The execution object used to use pointers to handle the virtual objects.
That is no longer necessary, so a simple `vtkm::exec::ColorTable` is
returned for execution objects. (Note that this `ColorTable` contains
pointers that are specific for the particular device.) This is a non-
backward compabible change. However, the only place (outside of the
`ColorTable` implementation itself) was a single worklet for converting
scalars to colors (`vtkm::worklet::colorconversion::TransferFunction`).
This is unlikely to affect anyone.
I also "fixed" some names in enum structs. There has been some
inconsistencies in VTK-m on whether items in an enum struct are
capitolized or camel case. We seem to moving toward camel case, so
deprecate some old names.
482266b44 Particle: explicitly make the copy ctor and assignment op
1037aa756 AssignerPartitionedDataSet: mark dtor as override
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2259
One of the features of `UnknownArrayHandle` is that it allows you to
query how many `Vec` components each value has without resolve the type
of the array. The functionality to implement this failed if you tried to
store an `ArrayHandle` that stored `Vec`-like objects with `Vec` sizes
that varied from value to value (i.e. an `ArrayHandleGroupVecVariable`).
Storing such an array in `UnknownArrayHandle` might not be the best
idea, but it should probably work. This change allows you to store such
an array. If you try to query the number of components, you will get 0.
e706880d7 Fix unnecessary deprecation warnings in visual studio
f7cc03107 Fix deprecated warnings
b27e4c7ea Ignore files for deprecated virtual classes for SourceInInstall test
284774cf4 Turn on "no virtual" option on some CI builds
cd08fd499 Add changelog for removing virtual methods
63ef84ed7 Optionally remove all use of ArrayHandleVirtual
e11f612ad Deprecate ArrayHandleVirtual
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2256
As we remove more and more virtual methods from VTK-m, I expect several
users will be interested in completely removing them from the build for
several reasons.
1. They may be compiling for hardware that does not support virtual
methods.
2. They may need to compile for CUDA but need shared libraries.
3. It should go a bit faster.
To enable this, a CMake option named `VTKm_NO_DEPRECATED_VIRTUAL` is
added. It defaults to `OFF`. But when it is `ON`, none of the code that
both uses virtuals and is deprecated will be built.
Currently, only `ArrayHandleVirtual` is deprecated, so the rest of the
virtual classes will still be built. As we move forward, more will be
removed until all virtual method functionality is removed.
386301719 Test UnknownArrayHandle behavior on special arrays
872da1f8e Suppress unnecessary deprecation warnings on VS
bb443bbc2 Replace the implementation of VariantArrayHandle
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !2255
`UnknownArrayHandle` has special behavior when putting in or getting out
an `ArrayHandleMultiplexer` or an `ArrayHandleCast`. When putting in
either of these, `UnknownArrayHandle` gets the actual array stored in
the multiplexer and cast so that you can later retrieve the base array.
Likewise, when you get the array out with `AsArrayHandle`, you can give
it an `ArrayHandleCast` or `ArrayHandleMultiplexer`, and you will get
the base array placed inside of it.
The Visual Studio compiler has an annoying habit where if you use a
templated class or method with a deprecated class as a template
parameter, you will get a deprecation warning where that class is used
in the templated thing. Thus, if you want to suppress the warning, you
have to supress every instance of the template, not just where the
template is declared.
This is annoying behavior that is thankfully not replicated in other
compilers.
The implementation of `VariantArrayHandle` has been changed to be a
relatively trivial subclass of `UnknownArrayHandle`.
The advantage of this change is twofold. First, it removes
`VariantArrayHandle`'s dependence on `ArrayHandleVirtual`, which gets us
much closer to deprecating that class. Second, it ensures that
`UnknownArrayHandle` is a reasonable replacement for
`VariantArrayHandle`, so we can move forward with replacing that.
C++11 introduced the `std::vector::data()` method. In addition to being
more syntatically pleasing, it should correctly handle the condition
when the `std::vector` is size 0 and therefore has no real pointer.
(Expressions like `&v[0]` are undefined when the `vector` is empty.)
The new method `ArrayHandle::DeepCopy` had the pattern such that the
data in the `this` array was pushed to the provided destination array.
However, this is the opposite pattern used in the equivalent method in
VTK, which takes the data from the provided array and copies it to
`this` array.
So, swap the pattern to match that of VTK. Also, make the method name
more descriptive by renaming it to `DeepCopyFrom`. Hopefully, users will
read that to mean given the `ArrayHandle`, you copy data from the other
provided `ArrayHandle`.
A recent modification to `ArrayCopy` created a fast path for when
copying arrays of the same type. This fast path just deep copies the
buffers. The issue was that the buffer copy was creating new buffers
instead of updating the existing buffers. The passed in `ArrayHandle`
would get updated to the new buffers, but any other `ArrayHandle`s
sharing those buffers would still have the old versions. That would lead
to unexpected errors in code like this.
```cpp
vtkm::cont::ArrayHandle<T> outArray1;
vtkm::cont::ArrayHandle<T> outArray2 = outArray1;
vtkm::cont::ArrayCopy(inArray, outArray2);
```
If `inArray` was a different type than `outArray2`, then the data for
both `outArray1` and `outArray2` would be updated (which is the expected
behavior for something considered a "pointer"). However, if `inArray`
was the same type as `outArray2`, then the fast path would change
`outArray2` but leave `outArray1` the same.
This behavior has been corrected so that, in this case, the data of
`outArray1` always follows that of `outArray2`.
The parallel sorts in the device adapter are documented to not be
stable. Up until recently, the sorts for all the supported devices
happened to be stable (or at least provided a stable sort where needed).
However, the recent Kokkos adapter provides a sort that is no stable.
This broke some tests that relied on stable sorts when they should not
have.
Fix the tests that relied on stable sort to check the results.
C++ was not resolving the overloads of `ArrayCopyImpl` as expected,
which was causing `ArrayCopy` to sometimes use a less efficient method
for copying.
Also fix an issue with `Buffer::DeepCopy` that caused a deadlock when
copying a buffer that was not actually allocated anywhere (as well as
failing to copy the metadata, which was probably the whole point).
The old style `ArrayHandle` stored most of its state, including the
data, in the `vtkm::cont::internal::Storage` object (templated to the
type of array). The new style of `ArrayHandle` stores the data itself in
`Buffer` objects, and recent changes to `Buffer` allow metadata to be
stored there, too.
These changes make it pretty unnecessary to hold any state at all in the
`Storage` object. This is good since the sharing of state from one type
of `ArrayHandle` to another (such as by transforming the data), can be
done by just sharing the `Buffer` objects.
To reinforce this behavior, the `Storage` object has been changed to
make it completely stateless. All the methods of `Storage` must be
marked as `static`.
While in the transition between two types of `ArrayHandle`
implementations, we need to declare when an `ArrayHandle` is implemented
with the new style. To consolidate, create a
`VTKM_ARRAY_HANDLE_NEW_STYLE` to override the default `ArrayHandle`
implementation with the `ArrayHandleNewStyle` implementation.
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
